[Bf-blender-cvs] [f56e59539f] cycles_split_kernel: Cycles: Remove remnants of tile splitting
Mai Lavelle
noreply at git.blender.org
Sat Feb 18 11:53:11 CET 2017
Commit: f56e59539fa24c6487de998cb147bf3f5c97bc2f
Author: Mai Lavelle
Date: Thu Feb 16 05:36:14 2017 -0500
Branches: cycles_split_kernel
https://developer.blender.org/rBf56e59539fa24c6487de998cb147bf3f5c97bc2f
Cycles: Remove remnants of tile splitting
===================================================================
M intern/cycles/device/device_split_kernel.cpp
M intern/cycles/device/opencl/opencl_split.cpp
M intern/cycles/kernel/kernels/opencl/kernel_data_init.cl
M intern/cycles/kernel/split/kernel_background_buffer_update.h
M intern/cycles/kernel/split/kernel_data_init.h
M intern/cycles/kernel/split/kernel_split_data.h
M intern/cycles/kernel/split/kernel_sum_all_radiance.h
M intern/cycles/render/buffers.cpp
M intern/cycles/render/buffers.h
===================================================================
diff --git a/intern/cycles/device/device_split_kernel.cpp b/intern/cycles/device/device_split_kernel.cpp
index 30315183f5..c5571ad36d 100644
--- a/intern/cycles/device/device_split_kernel.cpp
+++ b/intern/cycles/device/device_split_kernel.cpp
@@ -94,22 +94,6 @@ bool DeviceSplitKernel::path_trace(DeviceTask *task,
return false;
}
- /* TODO(mai): should be easy enough to remove these variables from tile */
- /* Buffer and rng_state offset calc. */
- {
- size_t offset_index = tile.offset + (tile.x + tile.y * tile.stride);
- size_t offset_x = offset_index % tile.stride;
- size_t offset_y = offset_index / tile.stride;
-
- tile.rng_state_offset_x = offset_x;
- tile.rng_state_offset_y = offset_y;
- tile.buffer_offset_x = offset_x;
- tile.buffer_offset_y = offset_y;
-
- tile.buffer_rng_state_stride = tile.stride;
- tile.stride = tile.w;
- }
-
/* Get local size */
size_t local_size[2];
{
diff --git a/intern/cycles/device/opencl/opencl_split.cpp b/intern/cycles/device/opencl/opencl_split.cpp
index 97b2ce495d..8a93fdc3c3 100644
--- a/intern/cycles/device/opencl/opencl_split.cpp
+++ b/intern/cycles/device/opencl/opencl_split.cpp
@@ -261,17 +261,11 @@ public:
rtile.h,
rtile.offset,
rtile.stride,
- rtile.rng_state_offset_x,
- rtile.rng_state_offset_y,
- rtile.buffer_rng_state_stride,
queue_index,
dQueue_size,
use_queues_flag,
work_pool_wgs,
rtile.num_samples,
- rtile.buffer_offset_x,
- rtile.buffer_offset_y,
- rtile.buffer_rng_state_stride,
rtile.buffer);
/* Enqueue ckPathTraceKernel_data_init kernel. */
diff --git a/intern/cycles/kernel/kernels/opencl/kernel_data_init.cl b/intern/cycles/kernel/kernels/opencl/kernel_data_init.cl
index c9315e3825..1e3c4fa28c 100644
--- a/intern/cycles/kernel/kernels/opencl/kernel_data_init.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel_data_init.cl
@@ -33,17 +33,11 @@ __kernel void kernel_ocl_path_trace_data_init(
int start_sample,
int end_sample,
int sx, int sy, int sw, int sh, int offset, int stride,
- int rng_state_offset_x,
- int rng_state_offset_y,
- int rng_state_stride,
ccl_global int *Queue_index, /* Tracks the number of elements in queues */
int queuesize, /* size (capacity) of the queue */
ccl_global char *use_queues_flag, /* flag to decide if scene-intersect kernel should use queues to fetch ray index */
ccl_global unsigned int *work_pool_wgs, /* Work pool for each work group */
unsigned int num_samples, /* Total number of samples per pixel */
- int buffer_offset_x,
- int buffer_offset_y,
- int buffer_stride,
ccl_global float *buffer)
{
kernel_data_init(kg,
@@ -59,16 +53,10 @@ __kernel void kernel_ocl_path_trace_data_init(
start_sample,
end_sample,
sx, sy, sw, sh, offset, stride,
- rng_state_offset_x,
- rng_state_offset_y,
- rng_state_stride,
Queue_index,
queuesize,
use_queues_flag,
work_pool_wgs,
num_samples,
- buffer_offset_x,
- buffer_offset_y,
- buffer_stride,
buffer);
}
diff --git a/intern/cycles/kernel/split/kernel_background_buffer_update.h b/intern/cycles/kernel/split/kernel_background_buffer_update.h
index 92755725a1..d6a9f3d09d 100644
--- a/intern/cycles/kernel/split/kernel_background_buffer_update.h
+++ b/intern/cycles/kernel/split/kernel_background_buffer_update.h
@@ -108,9 +108,6 @@ ccl_device void kernel_background_buffer_update(KernelGlobals *kg)
ccl_global uint *rng_state = kernel_split_params.rng_state;
int stride = kernel_split_params.stride;
- int rng_state_offset_x = kernel_split_params.rng_offset_x;
- int rng_state_offset_y = kernel_split_params.rng_offset_y;
- int rng_state_stride = kernel_split_params.rng_stride;
ccl_global char *ray_state = kernel_split_state.ray_state;
#ifdef __KERNEL_DEBUG__
@@ -143,7 +140,7 @@ ccl_device void kernel_background_buffer_update(KernelGlobals *kg)
my_sample_tile = 0;
initial_rng = rng_state;
- rng_state += (rng_state_offset_x + tile_x) + (rng_state_offset_y + tile_y) * rng_state_stride;
+ rng_state += tile_x + tile_y*kernel_split_params.stride;
per_sample_output_buffers += ((tile_x + (tile_y * stride)) + my_sample_tile) * kernel_data.film.pass_stride;
if(IS_STATE(ray_state, ray_index, RAY_HIT_BACKGROUND)) {
@@ -198,7 +195,7 @@ ccl_device void kernel_background_buffer_update(KernelGlobals *kg)
my_sample_tile = 0;
/* Remap rng_state according to the current work */
- rng_state = initial_rng + ((rng_state_offset_x + tile_x) + (rng_state_offset_y + tile_y) * rng_state_stride);
+ rng_state = initial_rng + tile_x + tile_y*kernel_split_params.stride;
/* Remap per_sample_output_buffers according to the current work */
per_sample_output_buffers = kernel_split_state.per_sample_output_buffers
+ ((tile_x + (tile_y * stride)) + my_sample_tile) * kernel_data.film.pass_stride;
diff --git a/intern/cycles/kernel/split/kernel_data_init.h b/intern/cycles/kernel/split/kernel_data_init.h
index 10f90c5196..3f38b0931d 100644
--- a/intern/cycles/kernel/split/kernel_data_init.h
+++ b/intern/cycles/kernel/split/kernel_data_init.h
@@ -72,17 +72,11 @@ void KERNEL_FUNCTION_FULL_NAME(data_init)(
int start_sample,
int end_sample,
int sx, int sy, int sw, int sh, int offset, int stride,
- int rng_state_offset_x,
- int rng_state_offset_y,
- int rng_state_stride,
ccl_global int *Queue_index, /* Tracks the number of elements in queues */
int queuesize, /* size (capacity) of the queue */
ccl_global char *use_queues_flag, /* flag to decide if scene-intersect kernel should use queues to fetch ray index */
ccl_global unsigned int *work_pools, /* Work pool for each work group */
- unsigned int num_samples, /* Total number of samples per pixel */
- int buffer_offset_x,
- int buffer_offset_y,
- int buffer_stride,
+ unsigned int num_samples,
ccl_global float *buffer)
{
#ifdef __KERNEL_OPENCL__
@@ -98,9 +92,6 @@ void KERNEL_FUNCTION_FULL_NAME(data_init)(
kernel_split_params.stride = stride;
kernel_split_params.rng_state = rng_state;
- kernel_split_params.rng_offset_x = rng_state_offset_x;
- kernel_split_params.rng_offset_y = rng_state_offset_y;
- kernel_split_params.rng_stride = rng_state_stride;
kernel_split_params.start_sample = start_sample;
kernel_split_params.end_sample = end_sample;
@@ -112,9 +103,6 @@ void KERNEL_FUNCTION_FULL_NAME(data_init)(
kernel_split_params.queue_size = queuesize;
kernel_split_params.use_queues_flag = use_queues_flag;
- kernel_split_params.buffer_offset_x = buffer_offset_x;
- kernel_split_params.buffer_offset_y = buffer_offset_y;
- kernel_split_params.buffer_stride = buffer_stride;
kernel_split_params.buffer = buffer;
split_data_init(&kernel_split_state, num_elements, split_data_buffer, ray_state);
@@ -185,7 +173,7 @@ void KERNEL_FUNCTION_FULL_NAME(data_init)(
ray_index);
kernel_split_state.work_array[ray_index] = work_index;
- rng_state += (rng_state_offset_x + tile_x) + (rng_state_offset_y + tile_y) * rng_state_stride;
+ rng_state += tile_x + tile_y*stride;
ccl_global float *per_sample_output_buffers = kernel_split_state.per_sample_output_buffers;
per_sample_output_buffers += ((tile_x + (tile_y * stride)) + (my_sample_tile)) * kernel_data.film.pass_stride;
diff --git a/intern/cycles/kernel/split/kernel_split_data.h b/intern/cycles/kernel/split/kernel_split_data.h
index 2ac84a49fe..7e88b6f516 100644
--- a/intern/cycles/kernel/split/kernel_split_data.h
+++ b/intern/cycles/kernel/split/kernel_split_data.h
@@ -31,9 +31,6 @@ typedef struct SplitParams {
int stride;
ccl_global uint *rng_state;
- int rng_offset_x;
- int rng_offset_y;
- int rng_stride;
int start_sample;
int end_sample;
@@ -45,9 +42,6 @@ typedef struct SplitParams {
int queue_size;
ccl_global char *use_queues_flag;
- int buffer_offset_x;
- int buffer_offset_y;
- int buffer_stride;
ccl_global float *buffer;
} SplitParams;
diff --git a/intern/cycles/kernel/split/kernel_sum_all_radiance.h b/intern/cycles/kernel/split/kernel_sum_all_radiance.h
index 958cbfc56a..485e144257 100644
--- a/intern/cycles/kernel/split/kernel_sum_all_radiance.h
+++ b/intern/cycles/kernel/split/kernel_sum_all_radiance.h
@@ -30,13 +30,10 @@ ccl_device void kernel_sum_all_radiance(KernelGlobals *kg)
int sw = kernel_split_params.w;
int sh = kernel_split_params.h;
int stride = kernel_split_params.stride;
- int buffer_offset_x = kernel_split_params.buffer_offset_x;
- int buffer_offset_y = kernel_split_params.buffer_offset_y;
- int buffer_stride = kernel_split_params.buffer_stride;
int start_sample = kernel_split_params.start_sample;
if(x < sw && y < sh) {
- buffer += ((buffer_offset_x + x) + (buffer_offset_y + y) * buffer_stride) * (kernel_data.film.pass_stride);
+ buffer += (x + y*stride) * (kernel_data.film.pass_stride);
ccl_global float *per_sample_output_buffer = kernel_split_state.per_sample_output_buffers;
per_sample_output_buffer += (x + y * stride) * (kernel_data.film.pass_stride);
diff --git a/intern/cycles/render/buffers.cpp b/intern/cycles/render/buffers.cpp
index 28481bf506..e3ef4bf13f 100644
--- a/intern/cycles/render/buffers.cpp
+++ b/intern/cycles/render/buffers.cpp
@@ -93,12 +93,6 @@ RenderTile::RenderTile()
rng_state = 0;
buffers = NULL;
-
- buffer_offset_x = 0;
- buffer_offset_y = 0;
- rng_state_offset_x = 0;
- rng_state_offset_y = 0;
- buffer_rng_state_stride = 0;
}
/* Render Buffers */
d
@@ Diff output truncated at 10240 characters. @@
More information about the Bf-blender-cvs
mailing list