Cycles: border rendering support, includes some refactoring in how pixels are

accessed on devices.

9 files changed, 41 insertions, 41 deletions

diff --git a/intern/cycles/kernel/kernel.cl b/intern/cycles/kernel/kernel.cl
index c00bc3fe957..90eb7a2513f 100644
--- a/intern/cycles/kernel/kernel.cl
+++ b/intern/cycles/kernel/kernel.cl
@@ -36,7 +36,7 @@ __kernel void kernel_ocl_path_trace(
 #include "kernel_textures.h"
 
 	int sample,
-	int sx, int sy, int sw, int sh)
+	int sx, int sy, int sw, int sh, int offset, int stride)
 {
 	KernelGlobals kglobals, *kg = &kglobals;
 
@@ -50,7 +50,7 @@ __kernel void kernel_ocl_path_trace(
 	int y = sy + get_global_id(1);
 
 	if(x < sx + sw && y < sy + sh)
-		kernel_path_trace(kg, buffer, rng_state, sample, x, y);
+		kernel_path_trace(kg, buffer, rng_state, sample, x, y, offset, stride);
 }
 
 __kernel void kernel_ocl_tonemap(
@@ -63,7 +63,7 @@ __kernel void kernel_ocl_tonemap(
 #include "kernel_textures.h"
 
 	int sample, int resolution,
-	int sx, int sy, int sw, int sh)
+	int sx, int sy, int sw, int sh, int offset, int stride)
 {
 	KernelGlobals kglobals, *kg = &kglobals;
 
@@ -77,7 +77,7 @@ __kernel void kernel_ocl_tonemap(
 	int y = sy + get_global_id(1);
 
 	if(x < sx + sw && y < sy + sh)
-		kernel_film_tonemap(kg, rgba, buffer, sample, resolution, x, y);
+		kernel_film_tonemap(kg, rgba, buffer, sample, resolution, x, y, offset, stride);
 }
 
 /*__kernel void kernel_ocl_displace(__global uint4 *input, __global float3 *offset, int sx)
diff --git a/intern/cycles/kernel/kernel.cpp b/intern/cycles/kernel/kernel.cpp
index 52a3852aa01..b4c3839dbd0 100644
--- a/intern/cycles/kernel/kernel.cpp
+++ b/intern/cycles/kernel/kernel.cpp
@@ -204,16 +204,16 @@ void kernel_tex_copy(KernelGlobals *kg, const char *name, device_ptr mem, size_t
 
 /* Path Tracing */
 
-void kernel_cpu_path_trace(KernelGlobals *kg, float4 *buffer, unsigned int *rng_state, int sample, int x, int y)
+void kernel_cpu_path_trace(KernelGlobals *kg, float4 *buffer, unsigned int *rng_state, int sample, int x, int y, int offset, int stride)
 {
-	kernel_path_trace(kg, buffer, rng_state, sample, x, y);
+	kernel_path_trace(kg, buffer, rng_state, sample, x, y, offset, stride);
 }
 
 /* Tonemapping */
 
-void kernel_cpu_tonemap(KernelGlobals *kg, uchar4 *rgba, float4 *buffer, int sample, int resolution, int x, int y)
+void kernel_cpu_tonemap(KernelGlobals *kg, uchar4 *rgba, float4 *buffer, int sample, int resolution, int x, int y, int offset, int stride)
 {
-	kernel_film_tonemap(kg, rgba, buffer, sample, resolution, x, y);
+	kernel_film_tonemap(kg, rgba, buffer, sample, resolution, x, y, offset, stride);
 }
 
 /* Displacement */
diff --git a/intern/cycles/kernel/kernel.cu b/intern/cycles/kernel/kernel.cu
index 75415a00b00..71fc7ac3197 100644
--- a/intern/cycles/kernel/kernel.cu
+++ b/intern/cycles/kernel/kernel.cu
@@ -26,22 +26,22 @@
 #include "kernel_path.h"
 #include "kernel_displace.h"
 
-extern "C" __global__ void kernel_cuda_path_trace(float4 *buffer, uint *rng_state, int sample, int sx, int sy, int sw, int sh)
+extern "C" __global__ void kernel_cuda_path_trace(float4 *buffer, uint *rng_state, int sample, int sx, int sy, int sw, int sh, int offset, int stride)
 {
 	int x = sx + blockDim.x*blockIdx.x + threadIdx.x;
 	int y = sy + blockDim.y*blockIdx.y + threadIdx.y;
 
 	if(x < sx + sw && y < sy + sh)
-		kernel_path_trace(NULL, buffer, rng_state, sample, x, y);
+		kernel_path_trace(NULL, buffer, rng_state, sample, x, y, offset, stride);
 }
 
-extern "C" __global__ void kernel_cuda_tonemap(uchar4 *rgba, float4 *buffer, int sample, int resolution, int sx, int sy, int sw, int sh)
+extern "C" __global__ void kernel_cuda_tonemap(uchar4 *rgba, float4 *buffer, int sample, int resolution, int sx, int sy, int sw, int sh, int offset, int stride)
 {
 	int x = sx + blockDim.x*blockIdx.x + threadIdx.x;
 	int y = sy + blockDim.y*blockIdx.y + threadIdx.y;
 
 	if(x < sx + sw && y < sy + sh)
-		kernel_film_tonemap(NULL, rgba, buffer, sample, resolution, x, y);
+		kernel_film_tonemap(NULL, rgba, buffer, sample, resolution, x, y, offset, stride);
 }
 
 extern "C" __global__ void kernel_cuda_displace(uint4 *input, float3 *offset, int sx)
diff --git a/intern/cycles/kernel/kernel.h b/intern/cycles/kernel/kernel.h
index 700ee49c5f2..78247504b39 100644
--- a/intern/cycles/kernel/kernel.h
+++ b/intern/cycles/kernel/kernel.h
@@ -36,13 +36,17 @@ bool kernel_osl_use(KernelGlobals *kg);
 void kernel_const_copy(KernelGlobals *kg, const char *name, void *host, size_t size);
 void kernel_tex_copy(KernelGlobals *kg, const char *name, device_ptr mem, size_t width, size_t height);
 
-void kernel_cpu_path_trace(KernelGlobals *kg, float4 *buffer, unsigned int *rng_state, int sample, int x, int y);
-void kernel_cpu_tonemap(KernelGlobals *kg, uchar4 *rgba, float4 *buffer, int sample, int resolution, int x, int y);
+void kernel_cpu_path_trace(KernelGlobals *kg, float4 *buffer, unsigned int *rng_state,
+	int sample, int x, int y, int offset, int stride);
+void kernel_cpu_tonemap(KernelGlobals *kg, uchar4 *rgba, float4 *buffer,
+	int sample, int resolution, int x, int y, int offset, int stride);
 void kernel_cpu_displace(KernelGlobals *kg, uint4 *input, float3 *offset, int i);
 
 #ifdef WITH_OPTIMIZED_KERNEL
-void kernel_cpu_optimized_path_trace(KernelGlobals *kg, float4 *buffer, unsigned int *rng_state, int sample, int x, int y);
-void kernel_cpu_optimized_tonemap(KernelGlobals *kg, uchar4 *rgba, float4 *buffer, int sample, int resolution, int x, int y);
+void kernel_cpu_optimized_path_trace(KernelGlobals *kg, float4 *buffer, unsigned int *rng_state,
+	int sample, int x, int y, int offset, int stride);
+void kernel_cpu_optimized_tonemap(KernelGlobals *kg, uchar4 *rgba, float4 *buffer,
+	int sample, int resolution, int x, int y, int offset, int stride);
 void kernel_cpu_optimized_displace(KernelGlobals *kg, uint4 *input, float3 *offset, int i);
 #endif
 
diff --git a/intern/cycles/kernel/kernel_film.h b/intern/cycles/kernel/kernel_film.h
index 4373701452e..cd8acc9647a 100644
--- a/intern/cycles/kernel/kernel_film.h
+++ b/intern/cycles/kernel/kernel_film.h
@@ -48,10 +48,9 @@ __device uchar4 film_float_to_byte(float4 color)
 	return result;
 }
 
-__device void kernel_film_tonemap(KernelGlobals *kg, __global uchar4 *rgba, __global float4 *buffer, int sample, int resolution, int x, int y)
+__device void kernel_film_tonemap(KernelGlobals *kg, __global uchar4 *rgba, __global float4 *buffer, int sample, int resolution, int x, int y, int offset, int stride)
 {
-	int w = kernel_data.cam.width;
-	int index = x + y*w;
+	int index = offset + x + y*stride;
 	float4 irradiance = buffer[index];
 
 	float4 float_result = film_map(kg, irradiance, sample);
diff --git a/intern/cycles/kernel/kernel_optimized.cpp b/intern/cycles/kernel/kernel_optimized.cpp
index 85a2b798a62..ea43e01ab58 100644
--- a/intern/cycles/kernel/kernel_optimized.cpp
+++ b/intern/cycles/kernel/kernel_optimized.cpp
@@ -35,16 +35,16 @@ CCL_NAMESPACE_BEGIN
 
 /* Path Tracing */
 
-void kernel_cpu_optimized_path_trace(KernelGlobals *kg, float4 *buffer, unsigned int *rng_state, int sample, int x, int y)
+void kernel_cpu_optimized_path_trace(KernelGlobals *kg, float4 *buffer, unsigned int *rng_state, int sample, int x, int y, int offset, int stride)
 {
-	kernel_path_trace(kg, buffer, rng_state, sample, x, y);
+	kernel_path_trace(kg, buffer, rng_state, sample, x, y, offset, stride);
 }
 
 /* Tonemapping */
 
-void kernel_cpu_optimized_tonemap(KernelGlobals *kg, uchar4 *rgba, float4 *buffer, int sample, int resolution, int x, int y)
+void kernel_cpu_optimized_tonemap(KernelGlobals *kg, uchar4 *rgba, float4 *buffer, int sample, int resolution, int x, int y, int offset, int stride)
 {
-	kernel_film_tonemap(kg, rgba, buffer, sample, resolution, x, y);
+	kernel_film_tonemap(kg, rgba, buffer, sample, resolution, x, y, offset, stride);
 }
 
 /* Displacement */
diff --git a/intern/cycles/kernel/kernel_path.h b/intern/cycles/kernel/kernel_path.h
index c609f6f13fe..05707f31352 100644
--- a/intern/cycles/kernel/kernel_path.h
+++ b/intern/cycles/kernel/kernel_path.h
@@ -34,7 +34,7 @@
 CCL_NAMESPACE_BEGIN
 
 #ifdef __MODIFY_TP__
-__device float3 path_terminate_modified_throughput(KernelGlobals *kg, __global float3 *buffer, int x, int y, int sample)
+__device float3 path_terminate_modified_throughput(KernelGlobals *kg, __global float3 *buffer, int x, int y, int offset, int stride, int sample)
 {
 	/* modify throughput to influence path termination probability, to avoid
 	   darker regions receiving fewer samples than lighter regions. also RGB
@@ -45,7 +45,7 @@ __device float3 path_terminate_modified_throughput(KernelGlobals *kg, __global f
 	const float minL = 0.1f;
 
 	if(sample >= minsample) {
-		float3 L = buffer[x + y*kernel_data.cam.width];
+		float3 L = buffer[offset + x + y*stride];
 		float3 Lmin = make_float3(minL, minL, minL);
 		float correct = (float)(sample+1)/(float)sample;
 
@@ -379,7 +379,7 @@ __device float4 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int sample, R
 	return make_float4(L.x, L.y, L.z, 1.0f - Ltransparent);
 }
 
-__device void kernel_path_trace(KernelGlobals *kg, __global float4 *buffer, __global uint *rng_state, int sample, int x, int y)
+__device void kernel_path_trace(KernelGlobals *kg, __global float4 *buffer, __global uint *rng_state, int sample, int x, int y, int offset, int stride)
 {
 	/* initialize random numbers */
 	RNG rng;
@@ -387,7 +387,7 @@ __device void kernel_path_trace(KernelGlobals *kg, __global float4 *buffer, __gl
 	float filter_u;
 	float filter_v;
 
-	path_rng_init(kg, rng_state, sample, &rng, x, y, &filter_u, &filter_v);
+	path_rng_init(kg, rng_state, sample, &rng, x, y, offset, stride, &filter_u, &filter_v);
 
 	/* sample camera ray */
 	Ray ray;
@@ -399,7 +399,7 @@ __device void kernel_path_trace(KernelGlobals *kg, __global float4 *buffer, __gl
 
 	/* integrate */
 #ifdef __MODIFY_TP__
-	float3 throughput = path_terminate_modified_throughput(kg, buffer, x, y, sample);
+	float3 throughput = path_terminate_modified_throughput(kg, buffer, x, y, offset, stride, sample);
 	float4 L = kernel_path_integrate(kg, &rng, sample, ray, throughput)/throughput;
 #else
 	float3 throughput = make_float3(1.0f, 1.0f, 1.0f);
@@ -407,14 +407,14 @@ __device void kernel_path_trace(KernelGlobals *kg, __global float4 *buffer, __gl
 #endif
 
 	/* accumulate result in output buffer */
-	int index = x + y*kernel_data.cam.width;
+	int index = offset + x + y*stride;
 
 	if(sample == 0)
 		buffer[index] = L;
 	else
 		buffer[index] += L;
 
-	path_rng_end(kg, rng_state, rng, x, y);
+	path_rng_end(kg, rng_state, rng, x, y, offset, stride);
 }
 
 CCL_NAMESPACE_END
diff --git a/intern/cycles/kernel/kernel_random.h b/intern/cycles/kernel/kernel_random.h
index ba97ab3e3b6..41301ebd3dc 100644
--- a/intern/cycles/kernel/kernel_random.h
+++ b/intern/cycles/kernel/kernel_random.h
@@ -123,7 +123,7 @@ __device_inline float path_rng(KernelGlobals *kg, RNG *rng, int sample, int dime
 #endif
 }
 
-__device_inline void path_rng_init(KernelGlobals *kg, __global uint *rng_state, int sample, RNG *rng, int x, int y, float *fx, float *fy)
+__device_inline void path_rng_init(KernelGlobals *kg, __global uint *rng_state, int sample, RNG *rng, int x, int y, int offset, int stride, float *fx, float *fy)
 {
 #ifdef __SOBOL_FULL_SCREEN__
 	uint px, py;
@@ -138,7 +138,7 @@ __device_inline void path_rng_init(KernelGlobals *kg, __global uint *rng_state,
 	*fx = size * (float)px * (1.0f/(float)0xFFFFFFFF) - x;
 	*fy = size * (float)py * (1.0f/(float)0xFFFFFFFF) - y;
 #else
-	*rng = rng_state[x + y*kernel_data.cam.width];
+	*rng = rng_state[offset + x + y*stride];
 
 	*rng ^= kernel_data.integrator.seed;
 
@@ -147,7 +147,7 @@ __device_inline void path_rng_init(KernelGlobals *kg, __global uint *rng_state,
 #endif
 }
 
-__device void path_rng_end(KernelGlobals *kg, __global uint *rng_state, RNG rng, int x, int y)
+__device void path_rng_end(KernelGlobals *kg, __global uint *rng_state, RNG rng, int x, int y, int offset, int stride)
 {
 	/* nothing to do */
 }
@@ -163,10 +163,10 @@ __device float path_rng(KernelGlobals *kg, RNG *rng, int sample, int dimension)
 	return (float)*rng * (1.0f/(float)0xFFFFFFFF);
 }
 
-__device void path_rng_init(KernelGlobals *kg, __global uint *rng_state, int sample, RNG *rng, int x, int y, float *fx, float *fy)
+__device void path_rng_init(KernelGlobals *kg, __global uint *rng_state, int sample, RNG *rng, int x, int y, int offset, int stride, float *fx, float *fy)
 {
 	/* load state */
-	*rng = rng_state[x + y*kernel_data.cam.width];
+	*rng = rng_state[offset + x + y*stride];
 
 	*rng ^= kernel_data.integrator.seed;
 
@@ -174,10 +174,10 @@ __device void path_rng_init(KernelGlobals *kg, __global uint *rng_state, int sam
 	*fy = path_rng(kg, rng, sample, PRNG_FILTER_V);
 }
 
-__device void path_rng_end(KernelGlobals *kg, __global uint *rng_state, RNG rng, int x, int y)
+__device void path_rng_end(KernelGlobals *kg, __global uint *rng_state, RNG rng, int x, int y, int offset, int stride)
 {
 	/* store state for next sample */
-	rng_state[x + y*kernel_data.cam.width] = rng;
+	rng_state[offset + x + y*stride] = rng;
 }
 
 #endif
diff --git a/intern/cycles/kernel/kernel_types.h b/intern/cycles/kernel/kernel_types.h
index d9bd645b16d..72ebfefbd90 100644
--- a/intern/cycles/kernel/kernel_types.h
+++ b/intern/cycles/kernel/kernel_types.h
@@ -300,10 +300,7 @@ typedef struct ShaderData {
 typedef struct KernelCamera {
 	/* type */
 	int ortho;
-	int pad;
-
-	/* size */
-	int width, height;
+	int pad1, pad2, pad3;
 
 	/* matrices */
 	Transform cameratoworld;