Cycles: border rendering support, includes some refactoring in how pixels are
authorBrecht Van Lommel <brechtvanlommel@pandora.be>
Tue, 20 Dec 2011 12:25:37 +0000 (12:25 +0000)
committerBrecht Van Lommel <brechtvanlommel@pandora.be>
Tue, 20 Dec 2011 12:25:37 +0000 (12:25 +0000)
accessed on devices.

25 files changed:
intern/cycles/app/cycles_test.cpp
intern/cycles/blender/addon/engine.py
intern/cycles/blender/blender_camera.cpp
intern/cycles/blender/blender_session.cpp
intern/cycles/blender/blender_sync.h
intern/cycles/device/device.h
intern/cycles/device/device_cpu.cpp
intern/cycles/device/device_cuda.cpp
intern/cycles/device/device_opencl.cpp
intern/cycles/kernel/kernel.cl
intern/cycles/kernel/kernel.cpp
intern/cycles/kernel/kernel.cu
intern/cycles/kernel/kernel.h
intern/cycles/kernel/kernel_film.h
intern/cycles/kernel/kernel_optimized.cpp
intern/cycles/kernel/kernel_path.h
intern/cycles/kernel/kernel_random.h
intern/cycles/kernel/kernel_types.h
intern/cycles/render/buffers.cpp
intern/cycles/render/buffers.h
intern/cycles/render/camera.cpp
intern/cycles/render/session.cpp
intern/cycles/render/session.h
intern/cycles/render/tile.cpp
intern/cycles/render/tile.h

index 27e53ded6db18775f89b5db3ef241994ebc0afa3..838167274048a0689678d74d26608274080a7018 100644 (file)
@@ -82,10 +82,21 @@ static void session_print_status()
        session_print(status);
 }
 
+static BufferParams session_buffer_params()
+{
+       BufferParams buffer_params;
+       buffer_params.width = options.width;
+       buffer_params.height = options.height;
+       buffer_params.full_width = options.width;
+       buffer_params.full_height = options.height;
+
+       return buffer_params;
+}
+
 static void session_init()
 {
        options.session = new Session(options.session_params);
-       options.session->reset(options.width, options.height, options.session_params.samples);
+       options.session->reset(session_buffer_params(), options.session_params.samples);
        options.session->scene = options.scene;
        
        if(options.session_params.background && !options.quiet)
@@ -151,7 +162,7 @@ static void display_info(Progress& progress)
 
 static void display()
 {
-       options.session->draw(options.width, options.height);
+       options.session->draw(session_buffer_params());
 
        display_info(options.session->progress);
 }
@@ -162,13 +173,13 @@ static void resize(int width, int height)
        options.height= height;
 
        if(options.session)
-               options.session->reset(options.width, options.height, options.session_params.samples);
+               options.session->reset(session_buffer_params(), options.session_params.samples);
 }
 
 void keyboard(unsigned char key)
 {
        if(key == 'r')
-               options.session->reset(options.width, options.height, options.session_params.samples);
+               options.session->reset(session_buffer_params(), options.session_params.samples);
        else if(key == 27) // escape
                options.session->progress.set_cancel("Cancelled");
 }
index 2fedd2c0afa060675fa73e928d1dcd046e783016..60b77b23f25c89ee5921cc57773d8ad8220d68f2 100644 (file)
@@ -62,11 +62,7 @@ def render(engine):
 
 def update(engine, data, scene):
     import bcycles
-    if scene.render.use_border:
-        engine.report({'ERROR'}, "Border rendering not supported yet")
-        free(engine)
-    else:
-        bcycles.sync(engine.session)
+    bcycles.sync(engine.session)
 
 
 def draw(engine, region, v3d, rv3d):
index 442a8f62bfd44f6ed20027c183e475032fc585e0..9777de14b1e98801c8572fa7cc596277bd3ef41a 100644 (file)
@@ -287,5 +287,29 @@ void BlenderSync::sync_view(BL::SpaceView3D b_v3d, BL::RegionView3D b_rv3d, int
        blender_camera_sync(scene->camera, &bcam, width, height);
 }
 
+BufferParams BlenderSync::get_buffer_params(BL::Scene b_scene, BL::RegionView3D b_rv3d, int width, int height)
+{
+       BufferParams params;
+
+       params.full_width = width;
+       params.full_height = height;
+
+       /* border render */
+       BL::RenderSettings r = b_scene.render();
+
+       if(!b_rv3d && r.use_border()) {
+               params.full_x = r.border_min_x()*width;
+               params.full_y = r.border_min_y()*height;
+               params.width = (int)(r.border_max_x()*width) - params.full_x;
+               params.height = (int)(r.border_max_y()*height) - params.full_y;
+       }
+       else {
+               params.width = width;
+               params.height = height;
+       }
+
+       return params;
+}
+
 CCL_NAMESPACE_END
 
index 4433b1e24f9149f65e6a0cdc5a1ce6bddb74d376..1803dd36beb43a650a548578be3c75137c19edb2 100644 (file)
@@ -100,7 +100,9 @@ void BlenderSession::create_session()
        session->set_pause(BlenderSync::get_session_pause(b_scene, background));
 
        /* start rendering */
-       session->reset(width, height, session_params.samples);
+       BufferParams buffer_params = BlenderSync::get_buffer_params(b_scene, b_rv3d, width, height);
+
+       session->reset(buffer_params, session_params.samples);
        session->start();
 }
 
@@ -135,7 +137,10 @@ void BlenderSession::write_render_result()
        if(!pixels)
                return;
 
-       struct RenderResult *rrp = RE_engine_begin_result((RenderEngine*)b_engine.ptr.data, 0, 0, width, height);
+       BufferParams buffer_params = BlenderSync::get_buffer_params(b_scene, b_rv3d, width, height);
+       int w = buffer_params.width, h = buffer_params.height;
+
+       struct RenderResult *rrp = RE_engine_begin_result((RenderEngine*)b_engine.ptr.data, 0, 0, w, h);
        PointerRNA rrptr;
        RNA_pointer_create(NULL, &RNA_RenderResult, rrp, &rrptr);
        BL::RenderResult rr(rrptr);
@@ -188,8 +193,10 @@ void BlenderSession::synchronize()
        session->scene->mutex.unlock();
 
        /* reset if needed */
-       if(scene->need_reset())
-               session->reset(width, height, session_params.samples);
+       if(scene->need_reset()) {
+               BufferParams buffer_params = BlenderSync::get_buffer_params(b_scene, b_rv3d, width, height);
+               session->reset(buffer_params, session_params.samples);
+       }
 }
 
 bool BlenderSession::draw(int w, int h)
@@ -225,7 +232,9 @@ bool BlenderSession::draw(int w, int h)
                /* reset if requested */
                if(reset) {
                        SessionParams session_params = BlenderSync::get_session_params(b_scene, background);
-                       session->reset(width, height, session_params.samples);
+                       BufferParams buffer_params = BlenderSync::get_buffer_params(b_scene, b_rv3d, width, height);
+
+                       session->reset(buffer_params, session_params.samples);
                }
        }
 
@@ -233,7 +242,9 @@ bool BlenderSession::draw(int w, int h)
        update_status_progress();
 
        /* draw */
-       return !session->draw(width, height);
+       BufferParams buffer_params = BlenderSync::get_buffer_params(b_scene, b_rv3d, width, height);
+
+       return !session->draw(buffer_params);
 }
 
 void BlenderSession::get_status(string& status, string& substatus)
index 83c7f70fd596876f07af39426189e9d312e0ac9b..824904cd81d2a64cbd2a87654a4a36ffc33a8c96 100644 (file)
@@ -62,6 +62,7 @@ public:
        static SceneParams get_scene_params(BL::Scene b_scene, bool background);
        static SessionParams get_session_params(BL::Scene b_scene, bool background);
        static bool get_session_pause(BL::Scene b_scene, bool background);
+       static BufferParams get_buffer_params(BL::Scene b_scene, BL::RegionView3D b_rv3d, int width, int height);
 
 private:
        /* sync */
index 5b87b11b6b87602b4167533432ab8ab018d20107..be6a3f144edd6ca75c9625133d44fdae3391d020 100644 (file)
@@ -60,6 +60,7 @@ public:
        device_ptr buffer;
        int sample;
        int resolution;
+       int offset, stride;
 
        device_ptr displace_input;
        device_ptr displace_offset;
index 990b7cb94b0a64d2909c73d89b3c3c7bb3a40fbb..a45a4fb69f627edd05900e355436715d39bccfb7 100644 (file)
@@ -162,7 +162,8 @@ public:
                if(system_cpu_support_optimized()) {
                        for(int y = task.y; y < task.y + task.h; y++) {
                                for(int x = task.x; x < task.x + task.w; x++)
-                                       kernel_cpu_optimized_path_trace(kg, (float4*)task.buffer, (unsigned int*)task.rng_state, task.sample, x, y);
+                                       kernel_cpu_optimized_path_trace(kg, (float4*)task.buffer, (unsigned int*)task.rng_state,
+                                               task.sample, x, y, task.offset, task.stride);
 
                                if(tasks.worker_cancel())
                                        break;
@@ -173,7 +174,8 @@ public:
                {
                        for(int y = task.y; y < task.y + task.h; y++) {
                                for(int x = task.x; x < task.x + task.w; x++)
-                                       kernel_cpu_path_trace(kg, (float4*)task.buffer, (unsigned int*)task.rng_state, task.sample, x, y);
+                                       kernel_cpu_path_trace(kg, (float4*)task.buffer, (unsigned int*)task.rng_state,
+                                               task.sample, x, y, task.offset, task.stride);
 
                                if(tasks.worker_cancel())
                                        break;
@@ -192,14 +194,16 @@ public:
                if(system_cpu_support_optimized()) {
                        for(int y = task.y; y < task.y + task.h; y++)
                                for(int x = task.x; x < task.x + task.w; x++)
-                                       kernel_cpu_optimized_tonemap(kg, (uchar4*)task.rgba, (float4*)task.buffer, task.sample, task.resolution, x, y);
+                                       kernel_cpu_optimized_tonemap(kg, (uchar4*)task.rgba, (float4*)task.buffer,
+                                               task.sample, task.resolution, x, y, task.offset, task.stride);
                }
                else
 #endif
                {
                        for(int y = task.y; y < task.y + task.h; y++)
                                for(int x = task.x; x < task.x + task.w; x++)
-                                       kernel_cpu_tonemap(kg, (uchar4*)task.rgba, (float4*)task.buffer, task.sample, task.resolution, x, y);
+                                       kernel_cpu_tonemap(kg, (uchar4*)task.rgba, (float4*)task.buffer,
+                                               task.sample, task.resolution, x, y, task.offset, task.stride);
                }
        }
 
index 177c90ba2dff3e4d72299bc99685c61d3150f2eb..dfa2fcb23222f27f88806fd5ae7c5e7c2ab59a3a 100644 (file)
@@ -520,6 +520,12 @@ public:
                cuda_assert(cuParamSeti(cuPathTrace, offset, task.h))
                offset += sizeof(task.h);
 
+               cuda_assert(cuParamSeti(cuPathTrace, offset, task.offset))
+               offset += sizeof(task.offset);
+
+               cuda_assert(cuParamSeti(cuPathTrace, offset, task.stride))
+               offset += sizeof(task.stride);
+
                cuda_assert(cuParamSetSize(cuPathTrace, offset))
 
                /* launch kernel: todo find optimal size, cache config for fermi */
@@ -581,6 +587,12 @@ public:
                cuda_assert(cuParamSeti(cuFilmConvert, offset, task.h))
                offset += sizeof(task.h);
 
+               cuda_assert(cuParamSeti(cuFilmConvert, offset, task.offset))
+               offset += sizeof(task.offset);
+
+               cuda_assert(cuParamSeti(cuFilmConvert, offset, task.stride))
+               offset += sizeof(task.stride);
+
                cuda_assert(cuParamSetSize(cuFilmConvert, offset))
 
                /* launch kernel: todo find optimal size, cache config for fermi */
index 8eaaebc6629c30f8c6e4f44cd5fc481e572ae9c6..3a1d3032d6e13bccfc319636d3ebf7db9eb62bf7 100644 (file)
@@ -191,7 +191,7 @@ public:
        {
                char version[256];
 
-               int major, minor, req_major = 1, req_minor = 1;
+               int major, minor, req_major = 1, req_minor = 0;
 
                clGetPlatformInfo(cpPlatform, CL_PLATFORM_VERSION, sizeof(version), &version, NULL);
 
@@ -541,6 +541,8 @@ public:
                cl_int d_w = task.w;
                cl_int d_h = task.h;
                cl_int d_sample = task.sample;
+               cl_int d_offset = task.offset;
+               cl_int d_stride = task.stride;
 
                /* sample arguments */
                int narg = 0;
@@ -559,6 +561,8 @@ public:
                ciErr |= clSetKernelArg(ckPathTraceKernel, narg++, sizeof(d_y), (void*)&d_y);
                ciErr |= clSetKernelArg(ckPathTraceKernel, narg++, sizeof(d_w), (void*)&d_w);
                ciErr |= clSetKernelArg(ckPathTraceKernel, narg++, sizeof(d_h), (void*)&d_h);
+               ciErr |= clSetKernelArg(ckPathTraceKernel, narg++, sizeof(d_offset), (void*)&d_offset);
+               ciErr |= clSetKernelArg(ckPathTraceKernel, narg++, sizeof(d_stride), (void*)&d_stride);
 
                opencl_assert(ciErr);
 
@@ -611,6 +615,8 @@ public:
                cl_int d_h = task.h;
                cl_int d_sample = task.sample;
                cl_int d_resolution = task.resolution;
+               cl_int d_offset = task.offset;
+               cl_int d_stride = task.stride;
 
                /* sample arguments */
                int narg = 0;
@@ -630,6 +636,8 @@ public:
                ciErr |= clSetKernelArg(ckFilmConvertKernel, narg++, sizeof(d_y), (void*)&d_y);
                ciErr |= clSetKernelArg(ckFilmConvertKernel, narg++, sizeof(d_w), (void*)&d_w);
                ciErr |= clSetKernelArg(ckFilmConvertKernel, narg++, sizeof(d_h), (void*)&d_h);
+               ciErr |= clSetKernelArg(ckFilmConvertKernel, narg++, sizeof(d_offset), (void*)&d_offset);
+               ciErr |= clSetKernelArg(ckFilmConvertKernel, narg++, sizeof(d_stride), (void*)&d_stride);
 
                opencl_assert(ciErr);
 
index c00bc3fe9579317130eaf0db0da164ba771864cf..90eb7a2513f2c34b30c48106bdb5a36498cb5dc8 100644 (file)
@@ -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)
index 52a3852aa01793fb46f51b3871917e0ccadaf21c..b4c3839dbd09824bb96327236ee4efa741d0ffef 100644 (file)
@@ -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 */
index 75415a00b000065111c6b61cac9e5c02ce5d5986..71fc7ac31971f9ea5c3d3bf4e1ef72efd674ac46 100644 (file)
 #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)
index 700ee49c5f2f98cc33dccf3683dd00a57bf4c585..78247504b39b36b791e3edbf88417a907ca8ce11 100644 (file)
@@ -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
 
index 4373701452ee080a4e85f666f526c9af8880b848..cd8acc9647abf76b5ef38b0586940b82b66f876a 100644 (file)
@@ -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);
index 85a2b798a628cc4beb7d162286bce3508eea4540..ea43e01ab58a65e9893c8920028b34ecbdb37f64 100644 (file)
@@ -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 */
index c609f6f13fe3722b13beebeb1f631243f3b83747..05707f31352e38fa778f53db6504dd522c4a1381 100644 (file)
@@ -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
index ba97ab3e3b65a8b50382c4b8ffaeacc6fc8b9a02..41301ebd3dcee79ad4238005cad39a5616c48537 100644 (file)
@@ -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
index d9bd645b16d35320f5b0841be4d29e6e3be7497f..72ebfefbd9082b97f4a9662dad18c18178593e2c 100644 (file)
@@ -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;
index acdddb475d0dbab492d162d9c1a07f2520a2302b..29141b25b59f5d70117089201b8eafdf797772d7 100644 (file)
@@ -36,8 +36,6 @@ CCL_NAMESPACE_BEGIN
 RenderBuffers::RenderBuffers(Device *device_)
 {
        device = device_;
-       width = 0;
-       height = 0;
 }
 
 RenderBuffers::~RenderBuffers()
@@ -58,24 +56,23 @@ void RenderBuffers::device_free()
        }
 }
 
-void RenderBuffers::reset(Device *device, int width_, int height_)
+void RenderBuffers::reset(Device *device, BufferParams& params_)
 {
-       width = width_;
-       height = height_;
+       params = params_;
 
        /* free existing buffers */
        device_free();
        
        /* allocate buffer */
-       buffer.resize(width, height);
+       buffer.resize(params.width, params.height);
        device->mem_alloc(buffer, MEM_READ_WRITE);
        device->mem_zero(buffer);
 
        /* allocate rng state */
-       rng_state.resize(width, height);
+       rng_state.resize(params.width, params.height);
 
-       uint *init_state = rng_state.resize(width, height);
-       int x, y;
+       uint *init_state = rng_state.resize(params.width, params.height);
+       int x, y, width = params.width, height = params.height;
        
        for(x=0; x<width; x++)
                for(y=0; y<height; y++)
@@ -92,11 +89,11 @@ float4 *RenderBuffers::copy_from_device(float exposure, int sample)
 
        device->mem_copy_from(buffer, 0, buffer.memory_size());
 
-       float4 *out = new float4[width*height];
+       float4 *out = new float4[params.width*params.height];
        float4 *in = (float4*)buffer.data_pointer;
        float scale = 1.0f/(float)sample;
        
-       for(int i = width*height - 1; i >= 0; i--) {
+       for(int i = params.width*params.height - 1; i >= 0; i--) {
                float4 rgba = in[i]*scale;
 
                rgba.x = rgba.x*exposure;
@@ -117,8 +114,6 @@ float4 *RenderBuffers::copy_from_device(float exposure, int sample)
 DisplayBuffer::DisplayBuffer(Device *device_)
 {
        device = device_;
-       width = 0;
-       height = 0;
        draw_width = 0;
        draw_height = 0;
        transparent = true; /* todo: determine from background */
@@ -137,28 +132,27 @@ void DisplayBuffer::device_free()
        }
 }
 
-void DisplayBuffer::reset(Device *device, int width_, int height_)
+void DisplayBuffer::reset(Device *device, BufferParams& params_)
 {
        draw_width = 0;
        draw_height = 0;
 
-       width = width_;
-       height = height_;
+       params = params_;
 
        /* free existing buffers */
        device_free();
 
        /* allocate display pixels */
-       rgba.resize(width, height);
+       rgba.resize(params.width, params.height);
        device->pixels_alloc(rgba);
 }
 
-void DisplayBuffer::draw_set(int width_, int height_)
+void DisplayBuffer::draw_set(int width, int height)
 {
-       assert(width_ <= width && height_ <= height);
+       assert(width <= params.width && height <= params.height);
 
-       draw_width = width_;
-       draw_height = height_;
+       draw_width = width;
+       draw_height = height;
 }
 
 void DisplayBuffer::draw_transparency_grid()
@@ -175,11 +169,11 @@ void DisplayBuffer::draw_transparency_grid()
        };
 
        glColor4ub(50, 50, 50, 255);
-       glRectf(0, 0, width, height);
+       glRectf(0, 0, params.width, params.height);
        glEnable(GL_POLYGON_STIPPLE);
        glColor4ub(55, 55, 55, 255);
        glPolygonStipple(checker_stipple_sml);
-       glRectf(0, 0, width, height);
+       glRectf(0, 0, params.width, params.height);
        glDisable(GL_POLYGON_STIPPLE);
 }
 
@@ -189,7 +183,7 @@ void DisplayBuffer::draw(Device *device)
                if(transparent)
                        draw_transparency_grid();
 
-               device->draw_pixels(rgba, 0, draw_width, draw_height, width, height, transparent);
+               device->draw_pixels(rgba, 0, draw_width, draw_height, params.width, params.height, transparent);
        }
 }
 
index d5eb8d7fa2feb92419d888463bc6acd488f614db..66bd03c889317f1c1f79b9035c8bf24d23d6df53 100644 (file)
@@ -30,12 +30,49 @@ CCL_NAMESPACE_BEGIN
 class Device;
 struct float4;
 
+/* Buffer Parameters
+   Size of render buffer and how it fits in the full image (border render). */
+
+class BufferParams {
+public:
+       /* width/height of the physical buffer */
+       int width;
+       int height;
+
+       /* offset into and width/height of the full buffer */
+       int full_x;
+       int full_y;
+       int full_width;
+       int full_height;
+
+       BufferParams()
+       {
+               width = 0;
+               height = 0;
+
+               full_x = 0;
+               full_y = 0;
+               full_width = 0;
+               full_height = 0;
+       }
+
+       bool modified(const BufferParams& params)
+       {
+               return !(full_x == params.full_x
+                       && full_y == params.full_y
+                       && width == params.width
+                       && height == params.height
+                       && full_width == params.full_width
+                       && full_height == params.full_height);
+       }
+};
+
 /* Render Buffers */
 
 class RenderBuffers {
 public:
-       /* buffer dimensions */
-       int width, height;
+       /* buffer parameters */
+       BufferParams params;
        /* float buffer */
        device_vector<float4> buffer;
        /* random number generator state */
@@ -46,7 +83,7 @@ public:
        RenderBuffers(Device *device);
        ~RenderBuffers();
 
-       void reset(Device *device, int width, int height);
+       void reset(Device *device, BufferParams& params);
        float4 *copy_from_device(float exposure, int sample);
 
 protected:
@@ -62,8 +99,8 @@ protected:
 
 class DisplayBuffer {
 public:
-       /* buffer dimensions */
-       int width, height;
+       /* buffer parameters */
+       BufferParams params;
        /* dimensions for how much of the buffer is actually ready for display.
           with progressive render we can be using only a subset of the buffer.
           if these are zero, it means nothing can be drawn yet */
@@ -78,7 +115,7 @@ public:
        DisplayBuffer(Device *device);
        ~DisplayBuffer();
 
-       void reset(Device *device, int width, int height);
+       void reset(Device *device, BufferParams& params);
        void write(Device *device, const string& filename);
 
        void draw_set(int width, int height);
index e88c0a388bc05e65b17d3bdcebee0313478ffb01..d5fca87491d3f2638cbcfe72278816306818de8e 100644 (file)
@@ -72,8 +72,9 @@ void Camera::update()
        if(!need_update)
                return;
        
+       /* ndc to raster */
        Transform screentocamera;
-       Transform ndctoraster = transform_scale((float)width, (float)height, 1.0f);
+       Transform ndctoraster = transform_scale(width, height, 1.0f);
 
        /* raster to screen */
        Transform screentoraster = ndctoraster *
@@ -148,10 +149,6 @@ void Camera::device_update(Device *device, DeviceScene *dscene)
        /* type */
        kcam->ortho = ortho;
 
-       /* size */
-       kcam->width = width;
-       kcam->height = height;
-
        /* store differentials */
        kcam->dx = dx;
        kcam->dy = dy;
index 42b4a2bb7e480de4ec35928aebc579f94dac46f2..26c4dbfbb7a7066e2819ff1b6b9970561a0f2409 100644 (file)
@@ -51,8 +51,6 @@ Session::Session(const SessionParams& params_)
        sample = 0;
 
        delayed_reset.do_reset = false;
-       delayed_reset.w = 0;
-       delayed_reset.h = 0;
        delayed_reset.samples = 0;
 
        display_outdated = false;
@@ -108,7 +106,7 @@ bool Session::ready_to_reset()
 
 /* GPU Session */
 
-void Session::reset_gpu(int w, int h, int samples)
+void Session::reset_gpu(BufferParams& buffer_params, int samples)
 {
        /* block for buffer acces and reset immediately. we can't do this
           in the thread, because we need to allocate an OpenGL buffer, and
@@ -119,7 +117,7 @@ void Session::reset_gpu(int w, int h, int samples)
        display_outdated = true;
        reset_time = time_dt();
 
-       reset_(w, h, samples);
+       reset_(buffer_params, samples);
 
        gpu_need_tonemap = false;
        gpu_need_tonemap_cond.notify_all();
@@ -127,7 +125,7 @@ void Session::reset_gpu(int w, int h, int samples)
        pause_cond.notify_all();
 }
 
-bool Session::draw_gpu(int w, int h)
+bool Session::draw_gpu(BufferParams& buffer_params)
 {
        /* block for buffer access */
        thread_scoped_lock display_lock(display->mutex);
@@ -136,7 +134,7 @@ bool Session::draw_gpu(int w, int h)
        if(gpu_draw_ready) {
                /* then verify the buffers have the expected size, so we don't
                   draw previous results in a resized window */
-               if(w == display->width && h == display->height) {
+               if(!buffer_params.modified(display->params)) {
                        /* for CUDA we need to do tonemapping still, since we can
                           only access GL buffers from the main thread */
                        if(gpu_need_tonemap) {
@@ -261,15 +259,14 @@ void Session::run_gpu()
 
 /* CPU Session */
 
-void Session::reset_cpu(int w, int h, int samples)
+void Session::reset_cpu(BufferParams& buffer_params, int samples)
 {
        thread_scoped_lock reset_lock(delayed_reset.mutex);
 
        display_outdated = true;
        reset_time = time_dt();
 
-       delayed_reset.w = w;
-       delayed_reset.h = h;
+       delayed_reset.params = buffer_params;
        delayed_reset.samples = samples;
        delayed_reset.do_reset = true;
        device->task_cancel();
@@ -277,7 +274,7 @@ void Session::reset_cpu(int w, int h, int samples)
        pause_cond.notify_all();
 }
 
-bool Session::draw_cpu(int w, int h)
+bool Session::draw_cpu(BufferParams& buffer_params)
 {
        thread_scoped_lock display_lock(display->mutex);
 
@@ -285,7 +282,7 @@ bool Session::draw_cpu(int w, int h)
        if(display->draw_ready()) {
                /* then verify the buffers have the expected size, so we don't
                   draw previous results in a resized window */
-               if(w == display->width && h == display->height) {
+               if(!buffer_params.modified(display->params)) {
                        display->draw(device);
 
                        if(display_outdated && (time_dt() - reset_time) > params.text_timeout)
@@ -306,7 +303,7 @@ void Session::run_cpu()
                thread_scoped_lock buffers_lock(buffers->mutex);
                thread_scoped_lock display_lock(display->mutex);
 
-               reset_(delayed_reset.w, delayed_reset.h, delayed_reset.samples);
+               reset_(delayed_reset.params, delayed_reset.samples);
                delayed_reset.do_reset = false;
        }
 
@@ -389,7 +386,7 @@ void Session::run_cpu()
                        if(delayed_reset.do_reset) {
                                /* reset rendering if request from main thread */
                                delayed_reset.do_reset = false;
-                               reset_(delayed_reset.w, delayed_reset.h, delayed_reset.samples);
+                               reset_(delayed_reset.params, delayed_reset.samples);
                        }
                        else if(need_tonemap) {
                                /* tonemap only if we do not reset, we don't we don't
@@ -438,23 +435,23 @@ void Session::run()
                progress.set_update();
 }
 
-bool Session::draw(int w, int h)
+bool Session::draw(BufferParams& buffer_params)
 {
        if(device_use_gl)
-               return draw_gpu(w, h);
+               return draw_gpu(buffer_params);
        else
-               return draw_cpu(w, h);
+               return draw_cpu(buffer_params);
 }
 
-void Session::reset_(int w, int h, int samples)
+void Session::reset_(BufferParams& buffer_params, int samples)
 {
-       if(w != buffers->width || h != buffers->height) {
+       if(buffer_params.modified(buffers->params)) {
                gpu_draw_ready = false;
-               buffers->reset(device, w, h);
-               display->reset(device, w, h);
+               buffers->reset(device, buffer_params);
+               display->reset(device, buffer_params);
        }
 
-       tile_manager.reset(w, h, samples);
+       tile_manager.reset(buffer_params, samples);
 
        start_time = time_dt();
        preview_time = 0.0;
@@ -462,12 +459,12 @@ void Session::reset_(int w, int h, int samples)
        sample = 0;
 }
 
-void Session::reset(int w, int h, int samples)
+void Session::reset(BufferParams& buffer_params, int samples)
 {
        if(device_use_gl)
-               reset_gpu(w, h, samples);
+               reset_gpu(buffer_params, samples);
        else
-               reset_cpu(w, h, samples);
+               reset_cpu(buffer_params, samples);
 }
 
 void Session::set_samples(int samples)
@@ -514,14 +511,18 @@ void Session::update_scene()
 
        progress.set_status("Updating Scene");
 
-       /* update camera if dimensions changed for progressive render */
+       /* update camera if dimensions changed for progressive render. the camera
+          knows nothing about progressive or cropped rendering, it just gets the
+          image dimensions passed in */
        Camera *cam = scene->camera;
-       int w = tile_manager.state.width;
-       int h = tile_manager.state.height;
-
-       if(cam->width != w || cam->height != h) {
-               cam->width = w;
-               cam->height = h;
+       float progressive_x = tile_manager.state.width/(float)tile_manager.params.width;
+       float progressive_y = tile_manager.state.height/(float)tile_manager.params.height;
+       int width = tile_manager.params.full_width*progressive_x;
+       int height = tile_manager.params.full_height*progressive_y;
+
+       if(width != cam->width || height != cam->height) {
+               cam->width = width;
+               cam->height = height;
                cam->tag_update();
        }
 
@@ -573,14 +574,16 @@ void Session::path_trace(Tile& tile)
        /* add path trace task */
        DeviceTask task(DeviceTask::PATH_TRACE);
 
-       task.x = tile.x;
-       task.y = tile.y;
+       task.x = tile_manager.state.full_x + tile.x;
+       task.y = tile_manager.state.full_y + tile.y;
        task.w = tile.w;
        task.h = tile.h;
        task.buffer = buffers->buffer.device_pointer;
        task.rng_state = buffers->rng_state.device_pointer;
        task.sample = tile_manager.state.sample;
        task.resolution = tile_manager.state.resolution;
+       task.offset = -(tile_manager.state.full_x + tile_manager.state.full_y*tile_manager.state.width);
+       task.stride = tile_manager.state.width;
 
        device->task_add(task);
 }
@@ -590,14 +593,16 @@ void Session::tonemap()
        /* add tonemap task */
        DeviceTask task(DeviceTask::TONEMAP);
 
-       task.x = 0;
-       task.y = 0;
+       task.x = tile_manager.state.full_x;
+       task.y = tile_manager.state.full_y;
        task.w = tile_manager.state.width;
        task.h = tile_manager.state.height;
        task.rgba = display->rgba.device_pointer;
        task.buffer = buffers->buffer.device_pointer;
        task.sample = tile_manager.state.sample;
        task.resolution = tile_manager.state.resolution;
+       task.offset = -(tile_manager.state.full_x + tile_manager.state.full_y*tile_manager.state.width);
+       task.stride = tile_manager.state.width;
 
        if(task.w > 0 && task.h > 0) {
                device->task_add(task);
index ce7f420096ae9e7637a65208afed73080ef29b5e..89979b8c451068adb5cd0569b8ab6d242ddf2633 100644 (file)
@@ -19,6 +19,7 @@
 #ifndef __SESSION_H__
 #define __SESSION_H__
 
+#include "buffers.h"
 #include "device.h"
 #include "tile.h"
 
@@ -27,6 +28,7 @@
 
 CCL_NAMESPACE_BEGIN
 
+class BufferParams;
 class Device;
 class DeviceScene;
 class DisplayBuffer;
@@ -106,11 +108,11 @@ public:
        ~Session();
 
        void start();
-       bool draw(int w, int h);
+       bool draw(BufferParams& params);
        void wait();
 
        bool ready_to_reset();
-       void reset(int w, int h, int samples);
+       void reset(BufferParams& params, int samples);
        void set_samples(int samples);
        void set_pause(bool pause);
 
@@ -118,7 +120,7 @@ protected:
        struct DelayedReset {
                thread_mutex mutex;
                bool do_reset;
-               int w, h;
+               BufferParams params;
                int samples;
        } delayed_reset;
 
@@ -129,15 +131,15 @@ protected:
 
        void tonemap();
        void path_trace(Tile& tile);
-       void reset_(int w, int h, int samples);
+       void reset_(BufferParams& params, int samples);
 
        void run_cpu();
-       bool draw_cpu(int w, int h);
-       void reset_cpu(int w, int h, int samples);
+       bool draw_cpu(BufferParams& params);
+       void reset_cpu(BufferParams& params, int samples);
 
        void run_gpu();
-       bool draw_gpu(int w, int h);
-       void reset_gpu(int w, int h, int samples);
+       bool draw_gpu(BufferParams& params);
+       void reset_gpu(BufferParams& params, int samples);
 
        TileManager tile_manager;
        bool device_use_gl;
index ba437e7487414d176e6144f76333c32717b2ff82..b118a7ba4781a276fe24ba480a72bdc688209a0c 100644 (file)
@@ -28,21 +28,21 @@ TileManager::TileManager(bool progressive_, int samples_, int tile_size_, int mi
        tile_size = tile_size_;
        min_size = min_size_;
 
-       reset(0, 0, 0);
+       BufferParams buffer_params;
+       reset(buffer_params, 0);
 }
 
 TileManager::~TileManager()
 {
 }
 
-void TileManager::reset(int width_, int height_, int samples_)
+void TileManager::reset(BufferParams& params_, int samples_)
 {
-       full_width = width_;
-       full_height = height_;
+       params = params_;
 
        start_resolution = 1;
 
-       int w = width_, h = height_;
+       int w = params.width, h = params.height;
 
        if(min_size != INT_MAX) {
                while(w*h > min_size*min_size) {
@@ -55,6 +55,8 @@ void TileManager::reset(int width_, int height_, int samples_)
 
        samples = samples_;
 
+       state.full_x = 0;
+       state.full_y = 0;
        state.width = 0;
        state.height = 0;
        state.sample = -1;
@@ -70,8 +72,8 @@ void TileManager::set_samples(int samples_)
 void TileManager::set_tiles()
 {
        int resolution = state.resolution;
-       int image_w = max(1, full_width/resolution);
-       int image_h = max(1, full_height/resolution);
+       int image_w = max(1, params.width/resolution);
+       int image_h = max(1, params.height/resolution);
        int tile_w = (image_w + tile_size - 1)/tile_size;
        int tile_h = (image_h + tile_size - 1)/tile_size;
        int sub_w = image_w/tile_w;
@@ -90,6 +92,8 @@ void TileManager::set_tiles()
                }
        }
 
+       state.full_x = params.full_x/resolution;
+       state.full_y = params.full_y/resolution;
        state.width = image_w;
        state.height = image_h;
 }
index 5cd16eb8afa9bed547a34277c1ea3e196d50fc15..76863d23498935abcf6b773925b626a2a04cd9c6 100644 (file)
@@ -21,6 +21,7 @@
 
 #include <limits.h>
 
+#include "buffers.h"
 #include "util_list.h"
 
 CCL_NAMESPACE_BEGIN
@@ -39,7 +40,10 @@ public:
 
 class TileManager {
 public:
+       BufferParams params;
        struct State {
+               int full_x;
+               int full_y;
                int width;
                int height;
                int sample;
@@ -50,7 +54,7 @@ public:
        TileManager(bool progressive, int samples, int tile_size, int min_size);
        ~TileManager();
 
-       void reset(int width, int height, int samples);
+       void reset(BufferParams& params, int samples);
        void set_samples(int samples);
        bool next();
        bool done();
@@ -63,8 +67,6 @@ protected:
        int tile_size;
        int min_size;
 
-       int full_width;
-       int full_height;
        int start_resolution;
 };