Merge branch 'master' into blender2.8

26 files changed, 358 insertions, 481 deletions

diff --git a/intern/cycles/blender/addon/properties.py b/intern/cycles/blender/addon/properties.py
index a50da3634f3..95a6e302cc5 100644
--- a/intern/cycles/blender/addon/properties.py
+++ b/intern/cycles/blender/addon/properties.py
@@ -1345,8 +1345,9 @@ class CyclesPreferences(bpy.types.AddonPreferences):
 
         cuda_devices = []
         opencl_devices = []
+        cpu_devices = []
         for device in device_list:
-            if not device[1] in {'CUDA', 'OPENCL'}:
+            if not device[1] in {'CUDA', 'OPENCL', 'CPU'}:
                 continue
 
             entry = None
@@ -1355,18 +1356,28 @@ class CyclesPreferences(bpy.types.AddonPreferences):
                 if dev.id == device[2] and dev.type == device[1]:
                     entry = dev
                     break
-            # Create new entry if no existing one was found
             if not entry:
+                # Create new entry if no existing one was found
                 entry = self.devices.add()
                 entry.id   = device[2]
                 entry.name = device[0]
                 entry.type = device[1]
+                entry.use  = entry.type != 'CPU'
+            elif entry.name != device[0]:
+                # Update name in case it changed
+                entry.name = device[0]
 
             # Sort entries into lists
             if entry.type == 'CUDA':
                 cuda_devices.append(entry)
             elif entry.type == 'OPENCL':
                 opencl_devices.append(entry)
+            else:
+                cpu_devices.append(entry)
+
+        cuda_devices.extend(cpu_devices)
+        opencl_devices.extend(cpu_devices)
+
         return cuda_devices, opencl_devices
 
 
diff --git a/intern/cycles/blender/addon/ui.py b/intern/cycles/blender/addon/ui.py
index 33b80b5695a..9c8855d0165 100644
--- a/intern/cycles/blender/addon/ui.py
+++ b/intern/cycles/blender/addon/ui.py
@@ -405,7 +405,12 @@ class CYCLES_RENDER_PT_performance(CyclesButtonsPanel, Panel):
         sub.prop(rd, "tile_x", text="X")
         sub.prop(rd, "tile_y", text="Y")
 
-        sub.prop(cscene, "use_progressive_refine")
+        subsub = sub.column()
+        subsub.active = not rd.use_save_buffers
+        for rl in rd.layers:
+            if rl.cycles.use_denoising:
+                subsub.active = False
+        subsub.prop(cscene, "use_progressive_refine")
 
         col = split.column()
 
@@ -595,7 +600,6 @@ class CYCLES_RENDER_PT_denoising(CyclesButtonsPanel, Panel):
         cscene = context.scene.cycles
         layout = self.layout
 
-        layout.active = not cscene.use_progressive_refine
         layout.prop(crl, "use_denoising", text="")
 
     def draw(self, context):
@@ -607,7 +611,7 @@ class CYCLES_RENDER_PT_denoising(CyclesButtonsPanel, Panel):
         rl = rd.layers.active
         crl = rl.cycles
 
-        layout.active = crl.use_denoising and not cscene.use_progressive_refine
+        layout.active = crl.use_denoising
 
         split = layout.split()
 
diff --git a/intern/cycles/blender/blender_session.cpp b/intern/cycles/blender/blender_session.cpp
index 75cc06ecd82..cac662d837e 100644
--- a/intern/cycles/blender/blender_session.cpp
+++ b/intern/cycles/blender/blender_session.cpp
@@ -311,11 +311,10 @@ static void end_render_result(BL::RenderEngine& b_engine,
 
 void BlenderSession::do_write_update_render_tile(RenderTile& rtile, bool do_update_only, bool highlight)
 {
-	BufferParams& params = rtile.buffers->params;
-	int x = params.full_x - session->tile_manager.params.full_x;
-	int y = params.full_y - session->tile_manager.params.full_y;
-	int w = params.width;
-	int h = params.height;
+	int x = rtile.x - session->tile_manager.params.full_x;
+	int y = rtile.y - session->tile_manager.params.full_y;
+	int w = rtile.w;
+	int h = rtile.h;
 
 	/* get render result */
 	BL::RenderResult b_rr = begin_render_result(b_engine, x, y, w, h, b_rlay_name.c_str(), b_rview_name.c_str());
@@ -407,7 +406,7 @@ void BlenderSession::render()
 		buffer_params.passes = passes;
 
 		PointerRNA crl = RNA_pointer_get(&b_layer_iter->ptr, "cycles");
-		bool use_denoising = !session_params.progressive_refine && get_boolean(crl, "use_denoising");
+		bool use_denoising = get_boolean(crl, "use_denoising");
 		buffer_params.denoising_data_pass = use_denoising;
 		session->tile_manager.schedule_denoising = use_denoising;
 		session->params.use_denoising = use_denoising;
@@ -665,10 +664,9 @@ void BlenderSession::do_write_update_render_result(BL::RenderResult& b_rr,
 	if(!buffers->copy_from_device())
 		return;
 
-	BufferParams& params = buffers->params;
 	float exposure = scene->film->exposure;
 
-	vector<float> pixels(params.width*params.height*4);
+	vector<float> pixels(rtile.w*rtile.h*4);
 
 	/* Adjust absolute sample number to the range. */
 	int sample = rtile.sample;
diff --git a/intern/cycles/blender/blender_sync.cpp b/intern/cycles/blender/blender_sync.cpp
index c1783045f66..3636c28ee99 100644
--- a/intern/cycles/blender/blender_sync.cpp
+++ b/intern/cycles/blender/blender_sync.cpp
@@ -558,8 +558,7 @@ array<Pass> BlenderSync::sync_render_passes(BL::RenderLayer& b_rlay,
 
 	PointerRNA crp = RNA_pointer_get(&b_srlay.ptr, "cycles");
 	if(get_boolean(crp, "denoising_store_passes") &&
-	   get_boolean(crp, "use_denoising") &&
-	   !session_params.progressive_refine) {
+	   get_boolean(crp, "use_denoising")) {
 		b_engine.add_pass("Denoising Normal",          3, "XYZ", b_srlay.name().c_str());
 		b_engine.add_pass("Denoising Normal Variance", 3, "XYZ", b_srlay.name().c_str());
 		b_engine.add_pass("Denoising Albedo",          3, "RGB", b_srlay.name().c_str());
@@ -660,6 +659,16 @@ SessionParams BlenderSync::get_session_params(BL::RenderEngine& b_engine,
 	/* feature set */
 	params.experimental = (get_enum(cscene, "feature_set") != 0);
 
+	/* threads */
+	BL::RenderSettings b_r = b_scene.render();
+	if(b_r.threads_mode() == BL::RenderSettings::threads_mode_FIXED)
+		params.threads = b_r.threads();
+	else
+		params.threads = 0;
+
+	/* Background */
+	params.background = background;
+
 	/* device type */
 	vector<DeviceInfo>& devices = Device::available_devices();
 	
@@ -688,12 +697,28 @@ SessionParams BlenderSync::get_session_params(BL::RenderEngine& b_engine,
 			}
 		}
 
-		int compute_device = get_enum(b_preferences, "compute_device_type");
+		enum ComputeDevice {
+			COMPUTE_DEVICE_CPU = 0,
+			COMPUTE_DEVICE_CUDA = 1,
+			COMPUTE_DEVICE_OPENCL = 2,
+			COMPUTE_DEVICE_NUM = 3,
+		};
+
+		ComputeDevice compute_device = (ComputeDevice)get_enum(b_preferences,
+		                                                       "compute_device_type",
+		                                                       COMPUTE_DEVICE_NUM,
+		                                                       COMPUTE_DEVICE_CPU);
 
-		if(compute_device != 0) {
+		if(compute_device != COMPUTE_DEVICE_CPU) {
 			vector<DeviceInfo> used_devices;
 			RNA_BEGIN(&b_preferences, device, "devices") {
-				if(get_enum(device, "type") == compute_device && get_boolean(device, "use")) {
+				ComputeDevice device_type = (ComputeDevice)get_enum(device,
+				                                                    "type",
+				                                                    COMPUTE_DEVICE_NUM,
+				                                                    COMPUTE_DEVICE_CPU);
+
+				if(get_boolean(device, "use") &&
+				   (device_type == compute_device || device_type == COMPUTE_DEVICE_CPU)) {
 					string id = get_string(device, "id");
 					foreach(DeviceInfo& info, devices) {
 						if(info.id == id) {
@@ -708,15 +733,14 @@ SessionParams BlenderSync::get_session_params(BL::RenderEngine& b_engine,
 				params.device = used_devices[0];
 			}
 			else if(used_devices.size() > 1) {
-				params.device = Device::get_multi_device(used_devices);
+				params.device = Device::get_multi_device(used_devices,
+				                                         params.threads,
+				                                         params.background);
 			}
 			/* Else keep using the CPU device that was set before. */
 		}
 	}
 
-	/* Background */
-	params.background = background;
-
 	/* samples */
 	int samples = get_int(cscene, "samples");
 	int aa_samples = get_int(cscene, "aa_samples");
@@ -776,20 +800,28 @@ SessionParams BlenderSync::get_session_params(BL::RenderEngine& b_engine,
 		params.tile_order = TILE_BOTTOM_TO_TOP;
 	}
 
+	/* other parameters */
 	params.start_resolution = get_int(cscene, "preview_start_resolution");
 	params.pixel_size = b_engine.get_preview_pixel_size(b_scene);
 
 	/* other parameters */
-	if(b_scene.render().threads_mode() == BL::RenderSettings::threads_mode_FIXED)
-		params.threads = b_scene.render().threads();
-	else
-		params.threads = 0;
-
 	params.cancel_timeout = (double)get_float(cscene, "debug_cancel_timeout");
 	params.reset_timeout = (double)get_float(cscene, "debug_reset_timeout");
 	params.text_timeout = (double)get_float(cscene, "debug_text_timeout");
 
-	params.progressive_refine = get_boolean(cscene, "use_progressive_refine");
+	/* progressive refine */
+	params.progressive_refine = get_boolean(cscene, "use_progressive_refine") &&
+	                            !b_r.use_save_buffers();
+
+	if(params.progressive_refine) {
+		BL::RenderSettings::layers_iterator b_rlay;
+		for(b_r.layers.begin(b_rlay); b_rlay != b_r.layers.end(); ++b_rlay) {
+			PointerRNA crl = RNA_pointer_get(&b_rlay->ptr, "cycles");
+			if(get_boolean(crl, "use_denoising")) {
+				params.progressive_refine = false;
+			}
+		}
+	}
 
 	if(background) {
 		if(params.progressive_refine)
diff --git a/intern/cycles/device/device.cpp b/intern/cycles/device/device.cpp
index 19c4bec55a8..5f01bd535d0 100644
--- a/intern/cycles/device/device.cpp
+++ b/intern/cycles/device/device.cpp
@@ -27,6 +27,7 @@
 #include "util/util_math.h"
 #include "util/util_opengl.h"
 #include "util/util_time.h"
+#include "util/util_system.h"
 #include "util/util_types.h"
 #include "util/util_vector.h"
 #include "util/util_string.h"
@@ -512,7 +513,7 @@ string Device::device_capabilities()
 	return capabilities;
 }
 
-DeviceInfo Device::get_multi_device(vector<DeviceInfo> subdevices)
+DeviceInfo Device::get_multi_device(const vector<DeviceInfo>& subdevices, int threads, bool background)
 {
 	assert(subdevices.size() > 1);
 
@@ -520,18 +521,38 @@ DeviceInfo Device::get_multi_device(vector<DeviceInfo> subdevices)
 	info.type = DEVICE_MULTI;
 	info.id = "MULTI";
 	info.description = "Multi Device";
-	info.multi_devices = subdevices;
 	info.num = 0;
 
 	info.has_bindless_textures = true;
 	info.has_volume_decoupled = true;
 	info.has_qbvh = true;
-	foreach(DeviceInfo &device, subdevices) {
-		assert(device.type == info.multi_devices[0].type);
-
+	foreach(const DeviceInfo &device, subdevices) {
 		info.has_bindless_textures &= device.has_bindless_textures;
 		info.has_volume_decoupled &= device.has_volume_decoupled;
 		info.has_qbvh &= device.has_qbvh;
+
+		if(device.type == DEVICE_CPU && subdevices.size() > 1) {
+			if(background) {
+				int orig_cpu_threads = (threads)? threads: system_cpu_thread_count();
+				int cpu_threads = max(orig_cpu_threads - (subdevices.size() - 1), 0);
+
+				if(cpu_threads >= 1) {
+					DeviceInfo cpu_device = device;
+					cpu_device.cpu_threads = cpu_threads;
+					info.multi_devices.push_back(cpu_device);
+				}
+
+				VLOG(1) << "CPU render threads reduced from "
+						<< orig_cpu_threads << " to " << cpu_threads
+						<< ", to dedicate to GPU.";
+			}
+			else {
+				VLOG(1) << "CPU render threads disabled for interactive render.";
+			}
+		}
+		else {
+			info.multi_devices.push_back(device);
+		}
 	}
 
 	return info;
diff --git a/intern/cycles/device/device.h b/intern/cycles/device/device.h
index 29803abd153..47332f52ace 100644
--- a/intern/cycles/device/device.h
+++ b/intern/cycles/device/device.h
@@ -58,6 +58,7 @@ public:
 	bool has_volume_decoupled;
 	bool has_qbvh;
 	bool use_split_kernel; /* Denotes if the device is going to run cycles using split-kernel */
+	int cpu_threads;
 	vector<DeviceInfo> multi_devices;
 
 	DeviceInfo()
@@ -65,6 +66,7 @@ public:
 		type = DEVICE_CPU;
 		id = "CPU";
 		num = 0;
+		cpu_threads = 0;
 		display_device = false;
 		advanced_shading = true;
 		has_bindless_textures = false;
@@ -359,7 +361,9 @@ public:
 	static vector<DeviceType>& available_types();
 	static vector<DeviceInfo>& available_devices();
 	static string device_capabilities();
-	static DeviceInfo get_multi_device(vector<DeviceInfo> subdevices);
+	static DeviceInfo get_multi_device(const vector<DeviceInfo>& subdevices,
+	                                   int threads,
+	                                   bool background);
 
 	/* Tag devices lists for update. */
 	static void tag_update();
diff --git a/intern/cycles/device/device_cpu.cpp b/intern/cycles/device/device_cpu.cpp
index a17caabc850..0ba00da16a6 100644
--- a/intern/cycles/device/device_cpu.cpp
+++ b/intern/cycles/device/device_cpu.cpp
@@ -207,8 +207,8 @@ public:
 	      KERNEL_NAME_EVAL(cpu_avx, name), \
 	      KERNEL_NAME_EVAL(cpu_avx2, name)
 
-	CPUDevice(DeviceInfo& info, Stats &stats, bool background)
-	: Device(info, stats, background),
+	CPUDevice(DeviceInfo& info_, Stats &stats_, bool background_)
+	: Device(info_, stats_, background_),
 #define REGISTER_KERNEL(name) name ## _kernel(KERNEL_FUNCTIONS(name))
 	  REGISTER_KERNEL(path_trace),
 	  REGISTER_KERNEL(convert_to_half_float),
@@ -229,6 +229,9 @@ public:
 	  REGISTER_KERNEL(data_init)
 #undef REGISTER_KERNEL
 	{
+		if(info.cpu_threads == 0) {
+			info.cpu_threads = TaskScheduler::num_threads();
+		}
 
 #ifdef WITH_OSL
 		kernel_globals.osl = &osl_globals;
@@ -237,7 +240,6 @@ public:
 		if(use_split_kernel) {
 			VLOG(1) << "Will be using split kernel.";
 		}
-
 		need_texture_info = false;
 
 #define REGISTER_SPLIT_KERNEL(name) split_kernels[#name] = KernelFunctions<void(*)(KernelGlobals*, KernelData*)>(KERNEL_FUNCTIONS(name))
@@ -271,7 +273,7 @@ public:
 
 	virtual bool show_samples() const
 	{
-		return (TaskScheduler::num_threads() == 1);
+		return (info.cpu_threads == 1);
 	}
 
 	void load_texture_info()
@@ -377,7 +379,7 @@ public:
 				texture_info.resize(flat_slot + 128);
 			}
 
-			TextureInfo& info = texture_info.get_data()[flat_slot];
+			TextureInfo& info = texture_info[flat_slot];
 			info.data = (uint64_t)mem.data_pointer;
 			info.cl_buffer = 0;
 			info.interpolation = interpolation;
@@ -826,9 +828,9 @@ public:
 	int get_split_task_count(DeviceTask& task)
 	{
 		if(task.type == DeviceTask::SHADER)
-			return task.get_subtask_count(TaskScheduler::num_threads(), 256);
+			return task.get_subtask_count(info.cpu_threads, 256);
 		else
-			return task.get_subtask_count(TaskScheduler::num_threads());
+			return task.get_subtask_count(info.cpu_threads);
 	}
 
 	void task_add(DeviceTask& task)
@@ -840,9 +842,9 @@ public:
 		list<DeviceTask> tasks;
 
 		if(task.type == DeviceTask::SHADER)
-			task.split(tasks, TaskScheduler::num_threads(), 256);
+			task.split(tasks, info.cpu_threads, 256);
 		else
-			task.split(tasks, TaskScheduler::num_threads());
+			task.split(tasks, info.cpu_threads);
 
 		foreach(DeviceTask& task, tasks)
 			task_pool.push(new CPUDeviceTask(this, task));
diff --git a/intern/cycles/device/device_cuda.cpp b/intern/cycles/device/device_cuda.cpp
index 44f9077188b..0f057e9966f 100644
--- a/intern/cycles/device/device_cuda.cpp
+++ b/intern/cycles/device/device_cuda.cpp
@@ -840,7 +840,7 @@ public:
 				}
 
 				/* Set Mapping and tag that we need to (re-)upload to device */
-				TextureInfo& info = texture_info.get_data()[flat_slot];
+				TextureInfo& info = texture_info[flat_slot];
 				info.data = (uint64_t)tex;
 				info.cl_buffer = 0;
 				info.interpolation = interpolation;
@@ -1932,9 +1932,10 @@ uint64_t CUDASplitKernel::state_buffer_size(device_memory& /*kg*/, device_memory
 	                           0, 0, (void**)&args, 0));
 
 	device->mem_copy_from(size_buffer, 0, 1, 1, sizeof(uint64_t));
+	size_t size = size_buffer[0];
 	device->mem_free(size_buffer);
 
-	return *size_buffer.get_data();
+	return size;
 }
 
 bool CUDASplitKernel::enqueue_split_kernel_data_init(const KernelDimensions& dim,
diff --git a/intern/cycles/device/device_memory.h b/intern/cycles/device/device_memory.h
index 20707ad04c9..eeeca61496e 100644
--- a/intern/cycles/device/device_memory.h
+++ b/intern/cycles/device/device_memory.h
@@ -270,31 +270,14 @@ public:
 		return &data[0];
 	}
 
-	T *copy(T *ptr, size_t width, size_t height = 0, size_t depth = 0)
+	void steal_data(array<T>& from)
 	{
-		T *mem = resize(width, height, depth);
-		if(mem != NULL) {
-			memcpy(mem, ptr, memory_size());
-		}
-		return mem;
-	}
-
-	void copy_at(T *ptr, size_t offset, size_t size)
-	{
-		if(size > 0) {
-			size_t mem_size = size*data_elements*datatype_size(data_type);
-			memcpy(&data[0] + offset, ptr, mem_size);
-		}
-	}
-
-	void reference(T *ptr, size_t width, size_t height = 0, size_t depth = 0)
-	{
-		data.clear();
-		data_size = width * ((height == 0)? 1: height) * ((depth == 0)? 1: depth);
-		data_pointer = (device_ptr)ptr;
-		data_width = width;
-		data_height = height;
-		data_depth = depth;
+		data.steal_data(from);
+		data_size = data.size();
+		data_pointer = (data_size)? (device_ptr)&data[0]: 0;
+		data_width = data_size;
+		data_height = 0;
+		data_depth = 0;
 	}
 
 	void clear()
@@ -318,6 +301,11 @@ public:
 		return &data[0];
 	}
 
+	T& operator[](size_t i)
+	{
+		return data[i];
+	}
+
 private:
 	array<T> data;
 };
diff --git a/intern/cycles/device/opencl/opencl_base.cpp b/intern/cycles/device/opencl/opencl_base.cpp
index 6747a8a83ac..48c32a9dc5c 100644
--- a/intern/cycles/device/opencl/opencl_base.cpp
+++ b/intern/cycles/device/opencl/opencl_base.cpp
@@ -494,20 +494,21 @@ void OpenCLDeviceBase::mem_free_sub_ptr(device_ptr device_pointer)
 void OpenCLDeviceBase::const_copy_to(const char *name, void *host, size_t size)
 {
 	ConstMemMap::iterator i = const_mem_map.find(name);
+	device_vector<uchar> *data;
 
 	if(i == const_mem_map.end()) {
-		device_vector<uchar> *data = new device_vector<uchar>();
-		data->copy((uchar*)host, size);
+		data = new device_vector<uchar>();
+		data->resize(size);
 
 		mem_alloc(name, *data, MEM_READ_ONLY);
-		i = const_mem_map.insert(ConstMemMap::value_type(name, data)).first;
+		const_mem_map.insert(ConstMemMap::value_type(name, data));
 	}
 	else {
-		device_vector<uchar> *data = i->second;
-		data->copy((uchar*)host, size);
+		data = i->second;
 	}
 
-	mem_copy_to(*i->second);
+	memcpy(data->get_data(), host, size);
+	mem_copy_to(*data);
 }
 
 void OpenCLDeviceBase::tex_alloc(const char *name,
diff --git a/intern/cycles/device/opencl/opencl_split.cpp b/intern/cycles/device/opencl/opencl_split.cpp
index b4e9419ebbd..920106f92d4 100644
--- a/intern/cycles/device/opencl/opencl_split.cpp
+++ b/intern/cycles/device/opencl/opencl_split.cpp
@@ -309,6 +309,7 @@ public:
 		device->opencl_assert_err(device->ciErr, "clEnqueueNDRangeKernel");
 
 		device->mem_copy_from(size_buffer, 0, 1, 1, sizeof(uint64_t));
+		size_t size = size_buffer[0];
 		device->mem_free(size_buffer);
 
 		if(device->ciErr != CL_SUCCESS) {
@@ -318,7 +319,7 @@ public:
 			return 0;
 		}
 
-		return *size_buffer.get_data();
+		return size;
 	}
 
 	virtual bool enqueue_split_kernel_data_init(const KernelDimensions& dim,
diff --git a/intern/cycles/device/opencl/opencl_util.cpp b/intern/cycles/device/opencl/opencl_util.cpp
index 7d5173a5f1d..459d512172f 100644
--- a/intern/cycles/device/opencl/opencl_util.cpp
+++ b/intern/cycles/device/opencl/opencl_util.cpp
@@ -1080,6 +1080,7 @@ cl_device_type OpenCLInfo::get_device_type(cl_device_id device_id)
 
 string OpenCLInfo::get_readable_device_name(cl_device_id device_id)
 {
+	string name = "";
 	char board_name[1024];
 	size_t length = 0;
 	if(clGetDeviceInfo(device_id,
@@ -1089,11 +1090,21 @@ string OpenCLInfo::get_readable_device_name(cl_device_id device_id)
 	                   &length) == CL_SUCCESS)
 	{
 		if(length != 0 && board_name[0] != '\0') {
-			return board_name;
+			name = board_name;
 		}
 	}
+
 	/* Fallback to standard device name API. */
-	return get_device_name(device_id);
+	if(name.empty()) {
+		name = get_device_name(device_id);
+	}
+
+	/* Distinguish from our native CPU device. */
+	if(get_device_type(device_id) & CL_DEVICE_TYPE_CPU) {
+		name += " (OpenCL)";
+	}
+
+	return name;
 }
 
 bool OpenCLInfo::get_driver_version(cl_device_id device_id,
diff --git a/intern/cycles/render/image.cpp b/intern/cycles/render/image.cpp
index bb94b9bb82a..c9fbd237010 100644
--- a/intern/cycles/render/image.cpp
+++ b/intern/cycles/render/image.cpp
@@ -364,6 +364,7 @@ int ImageManager::add_image(const string& filename,
 	img->extension = extension;
 	img->users = 1;
 	img->use_alpha = use_alpha;
+	img->mem = NULL;
 
 	images[type][slot] = img;
 
@@ -696,7 +697,6 @@ bool ImageManager::file_load_image(Image *img,
 }
 
 void ImageManager::device_load_image(Device *device,
-                                     DeviceScene *dscene,
                                      Scene *scene,
                                      ImageDataType type,
                                      int slot,
@@ -717,26 +717,27 @@ void ImageManager::device_load_image(Device *device,
 
 	/* Slot assignment */
 	int flat_slot = type_index_to_flattened_slot(slot, type);
-
 	string name = string_printf("__tex_image_%s_%03d", name_from_type(type).c_str(), flat_slot);
 
-	if(type == IMAGE_DATA_TYPE_FLOAT4) {
-		if(dscene->tex_float4_image[slot] == NULL)
-			dscene->tex_float4_image[slot] = new device_vector<float4>();
-		device_vector<float4>& tex_img = *dscene->tex_float4_image[slot];
+	/* Free previous texture in slot. */
+	if(img->mem) {
+		thread_scoped_lock device_lock(device_mutex);
+		device->tex_free(*img->mem);
+		delete img->mem;
+		img->mem = NULL;
+	}
 
-		if(tex_img.device_pointer) {
-			thread_scoped_lock device_lock(device_mutex);
-			device->tex_free(tex_img);
-		}
+	/* Create new texture. */
+	if(type == IMAGE_DATA_TYPE_FLOAT4) {
+		device_vector<float4> *tex_img = new device_vector<float4>();
 
 		if(!file_load_image<TypeDesc::FLOAT, float>(img,
 		                                            type,
 		                                            texture_limit,
-		                                            tex_img))
+		                                            *tex_img))
 		{
 			/* on failure to load, we set a 1x1 pixels pink image */
-			float *pixels = (float*)tex_img.resize(1, 1);
+			float *pixels = (float*)tex_img->resize(1, 1);
 
 			pixels[0] = TEX_IMAGE_MISSING_R;
 			pixels[1] = TEX_IMAGE_MISSING_G;
@@ -744,60 +745,34 @@ void ImageManager::device_load_image(Device *device,
 			pixels[3] = TEX_IMAGE_MISSING_A;
 		}
 
-		{
-			thread_scoped_lock device_lock(device_mutex);
-			device->tex_alloc(name.c_str(),
-			                  tex_img,
-			                  img->interpolation,
-			                  img->extension);
-		}
+		img->mem = tex_img;
 	}
 	else if(type == IMAGE_DATA_TYPE_FLOAT) {
-		if(dscene->tex_float_image[slot] == NULL)
-			dscene->tex_float_image[slot] = new device_vector<float>();
-		device_vector<float>& tex_img = *dscene->tex_float_image[slot];
-
-		if(tex_img.device_pointer) {
-			thread_scoped_lock device_lock(device_mutex);
-			device->tex_free(tex_img);
-		}
+		device_vector<float> *tex_img = new device_vector<float>();
 
 		if(!file_load_image<TypeDesc::FLOAT, float>(img,
 		                                            type,
 		                                            texture_limit,
-		                                            tex_img))
+		                                            *tex_img))
 		{
 			/* on failure to load, we set a 1x1 pixels pink image */
-			float *pixels = (float*)tex_img.resize(1, 1);
+			float *pixels = (float*)tex_img->resize(1, 1);
 
 			pixels[0] = TEX_IMAGE_MISSING_R;
 		}
 
-		{
-			thread_scoped_lock device_lock(device_mutex);
-			device->tex_alloc(name.c_str(),
-			                  tex_img,
-			                  img->interpolation,
-			                  img->extension);
-		}
+		img->mem = tex_img;
 	}
 	else if(type == IMAGE_DATA_TYPE_BYTE4) {
-		if(dscene->tex_byte4_image[slot] == NULL)
-			dscene->tex_byte4_image[slot] = new device_vector<uchar4>();
-		device_vector<uchar4>& tex_img = *dscene->tex_byte4_image[slot];
-
-		if(tex_img.device_pointer) {
-			thread_scoped_lock device_lock(device_mutex);
-			device->tex_free(tex_img);
-		}
+		device_vector<uchar4> *tex_img = new device_vector<uchar4>();
 
 		if(!file_load_image<TypeDesc::UINT8, uchar>(img,
 		                                            type,
 		                                            texture_limit,
-		                                            tex_img))
+		                                            *tex_img))
 		{
 			/* on failure to load, we set a 1x1 pixels pink image */
-			uchar *pixels = (uchar*)tex_img.resize(1, 1);
+			uchar *pixels = (uchar*)tex_img->resize(1, 1);
 
 			pixels[0] = (TEX_IMAGE_MISSING_R * 255);
 			pixels[1] = (TEX_IMAGE_MISSING_G * 255);
@@ -805,58 +780,32 @@ void ImageManager::device_load_image(Device *device,
 			pixels[3] = (TEX_IMAGE_MISSING_A * 255);
 		}
 
-		{
-			thread_scoped_lock device_lock(device_mutex);
-			device->tex_alloc(name.c_str(),
-			                  tex_img,
-			                  img->interpolation,
-			                  img->extension);
-		}
+		img->mem = tex_img;
 	}
-	else if(type == IMAGE_DATA_TYPE_BYTE){
-		if(dscene->tex_byte_image[slot] == NULL)
-			dscene->tex_byte_image[slot] = new device_vector<uchar>();
-		device_vector<uchar>& tex_img = *dscene->tex_byte_image[slot];
-
-		if(tex_img.device_pointer) {
-			thread_scoped_lock device_lock(device_mutex);
-			device->tex_free(tex_img);
-		}
+	else if(type == IMAGE_DATA_TYPE_BYTE) {
+		device_vector<uchar> *tex_img = new device_vector<uchar>();
 
 		if(!file_load_image<TypeDesc::UINT8, uchar>(img,
 		                                            type,
 		                                            texture_limit,
-		                                            tex_img)) {
+		                                            *tex_img)) {
 			/* on failure to load, we set a 1x1 pixels pink image */
-			uchar *pixels = (uchar*)tex_img.resize(1, 1);
+			uchar *pixels = (uchar*)tex_img->resize(1, 1);
 
 			pixels[0] = (TEX_IMAGE_MISSING_R * 255);
 		}
 
-		{
-			thread_scoped_lock device_lock(device_mutex);
-			device->tex_alloc(name.c_str(),
-			                  tex_img,
-			                  img->interpolation,
-			                  img->extension);
-		}
+		img->mem = tex_img;
 	}
-	else if(type == IMAGE_DATA_TYPE_HALF4){
-		if(dscene->tex_half4_image[slot] == NULL)
-			dscene->tex_half4_image[slot] = new device_vector<half4>();
-		device_vector<half4>& tex_img = *dscene->tex_half4_image[slot];
-
-		if(tex_img.device_pointer) {
-			thread_scoped_lock device_lock(device_mutex);
-			device->tex_free(tex_img);
-		}
+	else if(type == IMAGE_DATA_TYPE_HALF4) {
+		device_vector<half4> *tex_img = new device_vector<half4>();
 
 		if(!file_load_image<TypeDesc::HALF, half>(img,
 		                                          type,
 		                                          texture_limit,
-		                                          tex_img)) {
+		                                          *tex_img)) {
 			/* on failure to load, we set a 1x1 pixels pink image */
-			half *pixels = (half*)tex_img.resize(1, 1);
+			half *pixels = (half*)tex_img->resize(1, 1);
 
 			pixels[0] = TEX_IMAGE_MISSING_R;
 			pixels[1] = TEX_IMAGE_MISSING_G;
@@ -864,47 +813,38 @@ void ImageManager::device_load_image(Device *device,
 			pixels[3] = TEX_IMAGE_MISSING_A;
 		}
 
-		{
-			thread_scoped_lock device_lock(device_mutex);
-			device->tex_alloc(name.c_str(),
-			                  tex_img,
-			                  img->interpolation,
-			                  img->extension);
-		}
+		img->mem = tex_img;
 	}
-	else if(type == IMAGE_DATA_TYPE_HALF){
-		if(dscene->tex_half_image[slot] == NULL)
-			dscene->tex_half_image[slot] = new device_vector<half>();
-		device_vector<half>& tex_img = *dscene->tex_half_image[slot];
-
-		if(tex_img.device_pointer) {
-			thread_scoped_lock device_lock(device_mutex);
-			device->tex_free(tex_img);
-		}
+	else if(type == IMAGE_DATA_TYPE_HALF) {
+		device_vector<half> *tex_img = new device_vector<half>();
 
 		if(!file_load_image<TypeDesc::HALF, half>(img,
 		                                          type,
 		                                          texture_limit,
-		                                          tex_img)) {
+		                                          *tex_img)) {
 			/* on failure to load, we set a 1x1 pixels pink image */
-			half *pixels = (half*)tex_img.resize(1, 1);
+			half *pixels = (half*)tex_img->resize(1, 1);
 
 			pixels[0] = TEX_IMAGE_MISSING_R;
 		}
 
-		{
-			thread_scoped_lock device_lock(device_mutex);
-			device->tex_alloc(name.c_str(),
-			                  tex_img,
-			                  img->interpolation,
-			                  img->extension);
-		}
+		img->mem = tex_img;
 	}
 
+	/* Copy to device. */
+	if(img->mem) {
+		thread_scoped_lock device_lock(device_mutex);
+		device->tex_alloc(name.c_str(),
+						  *img->mem,
+						  img->interpolation,
+						  img->extension);
+	}
+
+
 	img->need_load = false;
 }
 
-void ImageManager::device_free_image(Device *device, DeviceScene *dscene, ImageDataType type, int slot)
+void ImageManager::device_free_image(Device *device, ImageDataType type, int slot)
 {
 	Image *img = images[type][slot];
 
@@ -915,105 +855,20 @@ void ImageManager::device_free_image(Device *device, DeviceScene *dscene, ImageD
 			((OSL::TextureSystem*)osl_texture_system)->invalidate(filename);
 #endif
 		}
-		else {
-			device_memory *tex_img = NULL;
-			switch(type) {
-				case IMAGE_DATA_TYPE_FLOAT4:
-					if(slot >= dscene->tex_float4_image.size()) {
-						break;
-					}
-					tex_img = dscene->tex_float4_image[slot];
-					dscene->tex_float4_image[slot] = NULL;
-					break;
-				case IMAGE_DATA_TYPE_BYTE4:
-					if(slot >= dscene->tex_byte4_image.size()) {
-						break;
-					}
-					tex_img = dscene->tex_byte4_image[slot];
-					dscene->tex_byte4_image[slot]= NULL;
-					break;
-				case IMAGE_DATA_TYPE_HALF4:
-					if(slot >= dscene->tex_half4_image.size()) {
-						break;
-					}
-					tex_img = dscene->tex_half4_image[slot];
-					dscene->tex_half4_image[slot]= NULL;
-					break;
-				case IMAGE_DATA_TYPE_FLOAT:
-					if(slot >= dscene->tex_float_image.size()) {
-						break;
-					}
-					tex_img = dscene->tex_float_image[slot];
-					dscene->tex_float_image[slot] = NULL;
-					break;
-				case IMAGE_DATA_TYPE_BYTE:
-					if(slot >= dscene->tex_byte_image.size()) {
-						break;
-					}
-					tex_img = dscene->tex_byte_image[slot];
-					dscene->tex_byte_image[slot]= NULL;
-					break;
-				case IMAGE_DATA_TYPE_HALF:
-					if(slot >= dscene->tex_half_image.size()) {
-						break;
-					}
-					tex_img = dscene->tex_half_image[slot];
-					dscene->tex_half_image[slot]= NULL;
-					break;
-				default:
-					assert(0);
-					tex_img = NULL;
-			}
-			if(tex_img) {
-				if(tex_img->device_pointer) {
-					thread_scoped_lock device_lock(device_mutex);
-					device->tex_free(*tex_img);
-				}
 
-				delete tex_img;
-			}
+		if(img->mem) {
+			thread_scoped_lock device_lock(device_mutex);
+			device->tex_free(*img->mem);
+			delete img->mem;
 		}
 
-		delete images[type][slot];
+		delete img;
 		images[type][slot] = NULL;
 		--tex_num_images[type];
 	}
 }
 
-void ImageManager::device_prepare_update(DeviceScene *dscene)
-{
-	for(int type = 0; type < IMAGE_DATA_NUM_TYPES; type++) {
-		switch(type) {
-			case IMAGE_DATA_TYPE_FLOAT4:
-				if(dscene->tex_float4_image.size() <= tex_num_images[IMAGE_DATA_TYPE_FLOAT4])
-					dscene->tex_float4_image.resize(tex_num_images[IMAGE_DATA_TYPE_FLOAT4]);
-				break;
-			case IMAGE_DATA_TYPE_BYTE4:
-				if(dscene->tex_byte4_image.size() <= tex_num_images[IMAGE_DATA_TYPE_BYTE4])
-					dscene->tex_byte4_image.resize(tex_num_images[IMAGE_DATA_TYPE_BYTE4]);
-				break;
-			case IMAGE_DATA_TYPE_HALF4:
-				if(dscene->tex_half4_image.size() <= tex_num_images[IMAGE_DATA_TYPE_HALF4])
-					dscene->tex_half4_image.resize(tex_num_images[IMAGE_DATA_TYPE_HALF4]);
-				break;
-			case IMAGE_DATA_TYPE_BYTE:
-				if(dscene->tex_byte_image.size() <= tex_num_images[IMAGE_DATA_TYPE_BYTE])
-					dscene->tex_byte_image.resize(tex_num_images[IMAGE_DATA_TYPE_BYTE]);
-				break;
-			case IMAGE_DATA_TYPE_FLOAT:
-				if(dscene->tex_float_image.size() <= tex_num_images[IMAGE_DATA_TYPE_FLOAT])
-					dscene->tex_float_image.resize(tex_num_images[IMAGE_DATA_TYPE_FLOAT]);
-				break;
-			case IMAGE_DATA_TYPE_HALF:
-				if(dscene->tex_half_image.size() <= tex_num_images[IMAGE_DATA_TYPE_HALF])
-					dscene->tex_half_image.resize(tex_num_images[IMAGE_DATA_TYPE_HALF]);
-				break;
-		}
-	}
-}
-
 void ImageManager::device_update(Device *device,
-                                 DeviceScene *dscene,
                                  Scene *scene,
                                  Progress& progress)
 {
@@ -1021,9 +876,6 @@ void ImageManager::device_update(Device *device,
 		return;
 	}
 
-	/* Make sure arrays are proper size. */
-	device_prepare_update(dscene);
-
 	TaskPool pool;
 	for(int type = 0; type < IMAGE_DATA_NUM_TYPES; type++) {
 		for(size_t slot = 0; slot < images[type].size(); slot++) {
@@ -1031,14 +883,13 @@ void ImageManager::device_update(Device *device,
 				continue;
 
 			if(images[type][slot]->users == 0) {
-				device_free_image(device, dscene, (ImageDataType)type, slot);
+				device_free_image(device, (ImageDataType)type, slot);
 			}
 			else if(images[type][slot]->need_load) {
 				if(!osl_texture_system || images[type][slot]->builtin_data)
 					pool.push(function_bind(&ImageManager::device_load_image,
 					                        this,
 					                        device,
-					                        dscene,
 					                        scene,
 					                        (ImageDataType)type,
 					                        slot,
@@ -1053,7 +904,6 @@ void ImageManager::device_update(Device *device,
 }
 
 void ImageManager::device_update_slot(Device *device,
-                                      DeviceScene *dscene,
                                       Scene *scene,
                                       int flat_slot,
                                       Progress *progress)
@@ -1065,12 +915,11 @@ void ImageManager::device_update_slot(Device *device,
 	assert(image != NULL);
 
 	if(image->users == 0) {
-		device_free_image(device, dscene, type, slot);
+		device_free_image(device, type, slot);
 	}
 	else if(image->need_load) {
 		if(!osl_texture_system || image->builtin_data)
 			device_load_image(device,
-			                  dscene,
 			                  scene,
 			                  type,
 			                  slot,
@@ -1078,31 +927,24 @@ void ImageManager::device_update_slot(Device *device,
 	}
 }
 
-void ImageManager::device_free_builtin(Device *device, DeviceScene *dscene)
+void ImageManager::device_free_builtin(Device *device)
 {
 	for(int type = 0; type < IMAGE_DATA_NUM_TYPES; type++) {
 		for(size_t slot = 0; slot < images[type].size(); slot++) {
 			if(images[type][slot] && images[type][slot]->builtin_data)
-				device_free_image(device, dscene, (ImageDataType)type, slot);
+				device_free_image(device, (ImageDataType)type, slot);
 		}
 	}
 }
 
-void ImageManager::device_free(Device *device, DeviceScene *dscene)
+void ImageManager::device_free(Device *device)
 {
 	for(int type = 0; type < IMAGE_DATA_NUM_TYPES; type++) {
 		for(size_t slot = 0; slot < images[type].size(); slot++) {
-			device_free_image(device, dscene, (ImageDataType)type, slot);
+			device_free_image(device, (ImageDataType)type, slot);
 		}
 		images[type].clear();
 	}
-
-	dscene->tex_float4_image.clear();
-	dscene->tex_byte4_image.clear();
-	dscene->tex_half4_image.clear();
-	dscene->tex_float_image.clear();
-	dscene->tex_byte_image.clear();
-	dscene->tex_half_image.clear();
 }
 
 CCL_NAMESPACE_END
diff --git a/intern/cycles/render/image.h b/intern/cycles/render/image.h
index c86d1cbedbf..cc7c8544bed 100644
--- a/intern/cycles/render/image.h
+++ b/intern/cycles/render/image.h
@@ -28,7 +28,6 @@
 CCL_NAMESPACE_BEGIN
 
 class Device;
-class DeviceScene;
 class Progress;
 class Scene;
 
@@ -62,18 +61,15 @@ public:
 	                                 bool& is_linear,
 	                                 bool& builtin_free_cache);
 
-	void device_prepare_update(DeviceScene *dscene);
 	void device_update(Device *device,
-	                   DeviceScene *dscene,
 	                   Scene *scene,
 	                   Progress& progress);
 	void device_update_slot(Device *device,
-	                        DeviceScene *dscene,
 	                        Scene *scene,
 	                        int flat_slot,
 	                        Progress *progress);
-	void device_free(Device *device, DeviceScene *dscene);
-	void device_free_builtin(Device *device, DeviceScene *dscene);
+	void device_free(Device *device);
+	void device_free_builtin(Device *device);
 
 	void set_osl_texture_system(void *texture_system);
 	bool set_animation_frame_update(int frame);
@@ -115,6 +111,8 @@ public:
 		InterpolationType interpolation;
 		ExtensionType extension;
 
+		device_memory *mem;
+
 		int users;
 	};
 
@@ -151,13 +149,11 @@ private:
 	string name_from_type(int type);
 
 	void device_load_image(Device *device,
-	                       DeviceScene *dscene,
 	                       Scene *scene,
 	                       ImageDataType type,
 	                       int slot,
 	                       Progress *progess);
 	void device_free_image(Device *device,
-	                       DeviceScene *dscene,
 	                       ImageDataType type,
 	                       int slot);
 };
diff --git a/intern/cycles/render/mesh.cpp b/intern/cycles/render/mesh.cpp
index 6470b3b1075..69c21fc3cb3 100644
--- a/intern/cycles/render/mesh.cpp
+++ b/intern/cycles/render/mesh.cpp
@@ -1127,14 +1127,12 @@ bool Mesh::is_instanced() const
 
 MeshManager::MeshManager()
 {
-	bvh = NULL;
 	need_update = true;
 	need_flags_update = true;
 }
 
 MeshManager::~MeshManager()
 {
-	delete bvh;
 }
 
 void MeshManager::update_osl_attributes(Device *device, Scene *scene, vector<AttributeRequestSet>& mesh_attributes)
@@ -1393,11 +1391,11 @@ static void update_attribute_element_size(Mesh *mesh,
 }
 
 static void update_attribute_element_offset(Mesh *mesh,
-                                            vector<float>& attr_float,
+                                            device_vector<float>& attr_float,
                                             size_t& attr_float_offset,
-                                            vector<float4>& attr_float3,
+                                            device_vector<float4>& attr_float3,
                                             size_t& attr_float3_offset,
-                                            vector<uchar4>& attr_uchar4,
+                                            device_vector<uchar4>& attr_uchar4,
                                             size_t& attr_uchar4_offset,
                                             Attribute *mattr,
                                             AttributePrimitive prim,
@@ -1425,7 +1423,7 @@ static void update_attribute_element_offset(Mesh *mesh,
 			uchar4 *data = mattr->data_uchar4();
 			offset = attr_uchar4_offset;
 
-			assert(attr_uchar4.capacity() >= offset + size);
+			assert(attr_uchar4.size() >= offset + size);
 			for(size_t k = 0; k < size; k++) {
 				attr_uchar4[offset+k] = data[k];
 			}
@@ -1435,7 +1433,7 @@ static void update_attribute_element_offset(Mesh *mesh,
 			float *data = mattr->data_float();
 			offset = attr_float_offset;
 
-			assert(attr_float.capacity() >= offset + size);
+			assert(attr_float.size() >= offset + size);
 			for(size_t k = 0; k < size; k++) {
 				attr_float[offset+k] = data[k];
 			}
@@ -1445,7 +1443,7 @@ static void update_attribute_element_offset(Mesh *mesh,
 			Transform *tfm = mattr->data_transform();
 			offset = attr_float3_offset;
 
-			assert(attr_float3.capacity() >= offset + size * 4);
+			assert(attr_float3.size() >= offset + size * 4);
 			for(size_t k = 0; k < size*4; k++) {
 				attr_float3[offset+k] = (&tfm->x)[k];
 			}
@@ -1455,7 +1453,7 @@ static void update_attribute_element_offset(Mesh *mesh,
 			float4 *data = mattr->data_float4();
 			offset = attr_float3_offset;
 
-			assert(attr_float3.capacity() >= offset + size);
+			assert(attr_float3.size() >= offset + size);
 			for(size_t k = 0; k < size; k++) {
 				attr_float3[offset+k] = data[k];
 			}
@@ -1556,9 +1554,9 @@ void MeshManager::device_update_attributes(Device *device, DeviceScene *dscene,
 		}
 	}
 
-	vector<float> attr_float(attr_float_size);
-	vector<float4> attr_float3(attr_float3_size);
-	vector<uchar4> attr_uchar4(attr_uchar4_size);
+	dscene->attributes_float.resize(attr_float_size);
+	dscene->attributes_float3.resize(attr_float3_size);
+	dscene->attributes_uchar4.resize(attr_uchar4_size);
 
 	size_t attr_float_offset = 0;
 	size_t attr_float3_offset = 0;
@@ -1577,27 +1575,27 @@ void MeshManager::device_update_attributes(Device *device, DeviceScene *dscene,
 			Attribute *subd_mattr = mesh->subd_attributes.find(req);
 
 			update_attribute_element_offset(mesh,
-			                                attr_float, attr_float_offset,
-			                                attr_float3, attr_float3_offset,
-			                                attr_uchar4, attr_uchar4_offset,
+			                                dscene->attributes_float, attr_float_offset,
+			                                dscene->attributes_float3, attr_float3_offset,
+			                                dscene->attributes_uchar4, attr_uchar4_offset,
 			                                triangle_mattr,
 			                                ATTR_PRIM_TRIANGLE,
 			                                req.triangle_type,
 			                                req.triangle_desc);
 
 			update_attribute_element_offset(mesh,
-			                                attr_float, attr_float_offset,
-			                                attr_float3, attr_float3_offset,
-			                                attr_uchar4, attr_uchar4_offset,
+			                                dscene->attributes_float, attr_float_offset,
+			                                dscene->attributes_float3, attr_float3_offset,
+			                                dscene->attributes_uchar4, attr_uchar4_offset,
 			                                curve_mattr,
 			                                ATTR_PRIM_CURVE,
 			                                req.curve_type,
 			                                req.curve_desc);
 
 			update_attribute_element_offset(mesh,
-			                                attr_float, attr_float_offset,
-			                                attr_float3, attr_float3_offset,
-			                                attr_uchar4, attr_uchar4_offset,
+			                                dscene->attributes_float, attr_float_offset,
+			                                dscene->attributes_float3, attr_float3_offset,
+			                                dscene->attributes_uchar4, attr_uchar4_offset,
 			                                subd_mattr,
 			                                ATTR_PRIM_SUBD,
 			                                req.subd_type,
@@ -1618,16 +1616,13 @@ void MeshManager::device_update_attributes(Device *device, DeviceScene *dscene,
 	/* copy to device */
 	progress.set_status("Updating Mesh", "Copying Attributes to device");
 
-	if(attr_float.size()) {
-		dscene->attributes_float.copy(&attr_float[0], attr_float.size());
+	if(dscene->attributes_float.size()) {
 		device->tex_alloc("__attributes_float", dscene->attributes_float);
 	}
-	if(attr_float3.size()) {
-		dscene->attributes_float3.copy(&attr_float3[0], attr_float3.size());
+	if(dscene->attributes_float3.size()) {
 		device->tex_alloc("__attributes_float3", dscene->attributes_float3);
 	}
-	if(attr_uchar4.size()) {
-		dscene->attributes_uchar4.copy(&attr_uchar4[0], attr_uchar4.size());
+	if(dscene->attributes_uchar4.size()) {
 		device->tex_alloc("__attributes_uchar4", dscene->attributes_uchar4);
 	}
 }
@@ -1725,10 +1720,9 @@ void MeshManager::device_update_mesh(Device *device,
 		}
 	}
 	else {
-		PackedBVH& pack = bvh->pack;
-		for(size_t i = 0; i < pack.prim_index.size(); ++i) {
-			if((pack.prim_type[i] & PRIMITIVE_ALL_TRIANGLE) != 0) {
-				tri_prim_index[pack.prim_index[i]] = pack.prim_tri_index[i];
+		for(size_t i = 0; i < dscene->prim_index.size(); ++i) {
+			if((dscene->prim_type[i] & PRIMITIVE_ALL_TRIANGLE) != 0) {
+				tri_prim_index[dscene->prim_index[i]] = dscene->prim_tri_index[i];
 			}
 		}
 	}
@@ -1832,11 +1826,13 @@ void MeshManager::device_update_bvh(Device *device, DeviceScene *dscene, Scene *
 	VLOG(1) << (bparams.use_qbvh ? "Using QBVH optimization structure"
 	                             : "Using regular BVH optimization structure");
 
-	delete bvh;
-	bvh = BVH::create(bparams, scene->objects);
+	BVH *bvh = BVH::create(bparams, scene->objects);
 	bvh->build(progress);
 
-	if(progress.get_cancel()) return;
+	if(progress.get_cancel()) {
+		delete bvh;
+		return;
+	}
 
 	/* copy to device */
 	progress.set_status("Updating Scene BVH", "Copying BVH to device");
@@ -1844,49 +1840,51 @@ void MeshManager::device_update_bvh(Device *device, DeviceScene *dscene, Scene *
 	PackedBVH& pack = bvh->pack;
 
 	if(pack.nodes.size()) {
-		dscene->bvh_nodes.reference((float4*)&pack.nodes[0], pack.nodes.size());
+		dscene->bvh_nodes.steal_data(pack.nodes);
 		device->tex_alloc("__bvh_nodes", dscene->bvh_nodes);
 	}
 	if(pack.leaf_nodes.size()) {
-		dscene->bvh_leaf_nodes.reference((float4*)&pack.leaf_nodes[0], pack.leaf_nodes.size());
+		dscene->bvh_leaf_nodes.steal_data(pack.leaf_nodes);
 		device->tex_alloc("__bvh_leaf_nodes", dscene->bvh_leaf_nodes);
 	}
 	if(pack.object_node.size()) {
-		dscene->object_node.reference((uint*)&pack.object_node[0], pack.object_node.size());
+		dscene->object_node.steal_data(pack.object_node);
 		device->tex_alloc("__object_node", dscene->object_node);
 	}
 	if(pack.prim_tri_index.size()) {
-		dscene->prim_tri_index.reference((uint*)&pack.prim_tri_index[0], pack.prim_tri_index.size());
+		dscene->prim_tri_index.steal_data(pack.prim_tri_index);
 		device->tex_alloc("__prim_tri_index", dscene->prim_tri_index);
 	}
 	if(pack.prim_tri_verts.size()) {
-		dscene->prim_tri_verts.reference((float4*)&pack.prim_tri_verts[0], pack.prim_tri_verts.size());
+		dscene->prim_tri_verts.steal_data(pack.prim_tri_verts);
 		device->tex_alloc("__prim_tri_verts", dscene->prim_tri_verts);
 	}
 	if(pack.prim_type.size()) {
-		dscene->prim_type.reference((uint*)&pack.prim_type[0], pack.prim_type.size());
+		dscene->prim_type.steal_data(pack.prim_type);
 		device->tex_alloc("__prim_type", dscene->prim_type);
 	}
 	if(pack.prim_visibility.size()) {
-		dscene->prim_visibility.reference((uint*)&pack.prim_visibility[0], pack.prim_visibility.size());
+		dscene->prim_visibility.steal_data(pack.prim_visibility);
 		device->tex_alloc("__prim_visibility", dscene->prim_visibility);
 	}
 	if(pack.prim_index.size()) {
-		dscene->prim_index.reference((uint*)&pack.prim_index[0], pack.prim_index.size());
+		dscene->prim_index.steal_data(pack.prim_index);
 		device->tex_alloc("__prim_index", dscene->prim_index);
 	}
 	if(pack.prim_object.size()) {
-		dscene->prim_object.reference((uint*)&pack.prim_object[0], pack.prim_object.size());
+		dscene->prim_object.steal_data(pack.prim_object);
 		device->tex_alloc("__prim_object", dscene->prim_object);
 	}
 	if(pack.prim_time.size()) {
-		dscene->prim_time.reference((float2*)&pack.prim_time[0], pack.prim_time.size());
+		dscene->prim_time.steal_data(pack.prim_time);
 		device->tex_alloc("__prim_time", dscene->prim_time);
 	}
 
 	dscene->data.bvh.root = pack.root_index;
 	dscene->data.bvh.use_qbvh = bparams.use_qbvh;
 	dscene->data.bvh.use_bvh_steps = (scene->params.num_bvh_time_steps != 0);
+
+	delete bvh;
 }
 
 void MeshManager::device_update_flags(Device * /*device*/,
@@ -1913,7 +1911,6 @@ void MeshManager::device_update_flags(Device * /*device*/,
 }
 
 void MeshManager::device_update_displacement_images(Device *device,
-                                                    DeviceScene *dscene,
                                                     Scene *scene,
                                                     Progress& progress)
 {
@@ -1941,12 +1938,10 @@ void MeshManager::device_update_displacement_images(Device *device,
 			}
 		}
 	}
-	image_manager->device_prepare_update(dscene);
 	foreach(int slot, bump_images) {
 		pool.push(function_bind(&ImageManager::device_update_slot,
 		                        image_manager,
 		                        device,
-		                        dscene,
 		                        scene,
 		                        slot,
 		                        &progress));
@@ -2029,7 +2024,7 @@ void MeshManager::device_update(Device *device, DeviceScene *dscene, Scene *scen
 	}
 	if(true_displacement_used) {
 		VLOG(1) << "Updating images used for true displacement.";
-		device_update_displacement_images(device, dscene, scene, progress);
+		device_update_displacement_images(device, scene, progress);
 		old_need_object_flags_update = scene->object_manager->need_flags_update;
 		scene->object_manager->device_update_flags(device,
 		                                           dscene,
@@ -2171,6 +2166,7 @@ void MeshManager::device_free(Device *device, DeviceScene *dscene)
 	device->tex_free(dscene->attributes_uchar4);
 
 	dscene->bvh_nodes.clear();
+	dscene->bvh_leaf_nodes.clear();
 	dscene->object_node.clear();
 	dscene->prim_tri_verts.clear();
 	dscene->prim_tri_index.clear();
diff --git a/intern/cycles/render/mesh.h b/intern/cycles/render/mesh.h
index ed7cb881e91..30f5e9063e6 100644
--- a/intern/cycles/render/mesh.h
+++ b/intern/cycles/render/mesh.h
@@ -321,8 +321,6 @@ public:
 
 class MeshManager {
 public:
-	BVH *bvh;
-
 	bool need_update;
 	bool need_flags_update;
 
@@ -368,7 +366,6 @@ protected:
 	                       Progress& progress);
 
 	void device_update_displacement_images(Device *device,
-	                                       DeviceScene *dscene,
 	                                       Scene *scene,
 	                                       Progress& progress);
 };
diff --git a/intern/cycles/render/scene.cpp b/intern/cycles/render/scene.cpp
index c59a5d97df5..cf89385a33d 100644
--- a/intern/cycles/render/scene.cpp
+++ b/intern/cycles/render/scene.cpp
@@ -107,9 +107,9 @@ void Scene::free_memory(bool final)
 		bake_manager->device_free(device, &dscene);
 
 		if(!params.persistent_data || final)
-			image_manager->device_free(device, &dscene);
+			image_manager->device_free(device);
 		else
-			image_manager->device_free_builtin(device, &dscene);
+			image_manager->device_free_builtin(device);
 
 		lookup_tables->device_free(device, &dscene);
 	}
@@ -185,7 +185,7 @@ void Scene::device_update(Device *device_, Progress& progress)
 	if(progress.get_cancel() || device->have_error()) return;
 
 	progress.set_status("Updating Images");
-	image_manager->device_update(device, &dscene, this, progress);
+	image_manager->device_update(device, this, progress);
 
 	if(progress.get_cancel() || device->have_error()) return;
 
diff --git a/intern/cycles/render/scene.h b/intern/cycles/render/scene.h
index a1966afd23b..d4ec7d90ff5 100644
--- a/intern/cycles/render/scene.h
+++ b/intern/cycles/render/scene.h
@@ -60,15 +60,15 @@ class BakeData;
 class DeviceScene {
 public:
 	/* BVH */
-	device_vector<float4> bvh_nodes;
-	device_vector<float4> bvh_leaf_nodes;
-	device_vector<uint> object_node;
+	device_vector<int4> bvh_nodes;
+	device_vector<int4> bvh_leaf_nodes;
+	device_vector<int> object_node;
 	device_vector<uint> prim_tri_index;
 	device_vector<float4> prim_tri_verts;
-	device_vector<uint> prim_type;
+	device_vector<int> prim_type;
 	device_vector<uint> prim_visibility;
-	device_vector<uint> prim_index;
-	device_vector<uint> prim_object;
+	device_vector<int> prim_index;
+	device_vector<int> prim_object;
 	device_vector<float2> prim_time;
 
 	/* mesh */
@@ -103,7 +103,7 @@ public:
 	device_vector<float4> particles;
 
 	/* shaders */
-	device_vector<uint4> svm_nodes;
+	device_vector<int4> svm_nodes;
 	device_vector<uint> shader_flag;
 	device_vector<uint> object_flag;
 
@@ -113,14 +113,6 @@ public:
 	/* integrator */
 	device_vector<uint> sobol_directions;
 
-	/* cpu images */
-	vector<device_vector<float4>* > tex_float4_image;
-	vector<device_vector<uchar4>* > tex_byte4_image;
-	vector<device_vector<half4>* > tex_half4_image;
-	vector<device_vector<float>* > tex_float_image;
-	vector<device_vector<uchar>* > tex_byte_image;
-	vector<device_vector<half>* > tex_half_image;
-
 	KernelData data;
 };
 
diff --git a/intern/cycles/render/session.cpp b/intern/cycles/render/session.cpp
index f1ff6b49b71..4642dcfa9a1 100644
--- a/intern/cycles/render/session.cpp
+++ b/intern/cycles/render/session.cpp
@@ -114,9 +114,7 @@ Session::~Session()
 	}
 
 	/* clean up */
-	foreach(RenderTile &rtile, render_tiles)
-		delete rtile.buffers;
-	tile_manager.free_device();
+	tile_manager.device_free();
 
 	delete buffers;
 	delete display;
@@ -380,19 +378,17 @@ bool Session::acquire_tile(Device *tile_device, RenderTile& rtile)
 
 	/* in case of a permanent buffer, return it, otherwise we will allocate
 	 * a new temporary buffer */
-	if(!(params.background && params.output_path.empty())) {
+	if(buffers) {
 		tile_manager.state.buffer.get_offset_stride(rtile.offset, rtile.stride);
 
 		rtile.buffer = buffers->buffer.device_pointer;
 		rtile.buffers = buffers;
-		tile->buffers = buffers;
 
 		device->map_tile(tile_device, rtile);
 
 		return true;
 	}
 
-	bool store_rtile = false;
 	if(tile->buffers == NULL) {
 		/* fill buffer parameters */
 		BufferParams buffer_params = tile_manager.params;
@@ -402,52 +398,15 @@ bool Session::acquire_tile(Device *tile_device, RenderTile& rtile)
 		buffer_params.height = rtile.h;
 
 		/* allocate buffers */
-		if(params.progressive_refine) {
-			tile_lock.lock();
-
-			if(render_tiles.size() == 0) {
-				RenderTile nulltile;
-				nulltile.buffers = NULL;
-				render_tiles.resize(tile_manager.state.num_tiles, nulltile);
-			}
-
-			/* In certain circumstances number of tiles in the tile manager could
-			 * be changed. This is not supported by the progressive refine feature.
-			 */
-			assert(render_tiles.size() == tile_manager.state.num_tiles);
-
-			RenderTile &stored_rtile = render_tiles[tile->index];
-			if(stored_rtile.buffers == NULL) {
-				tile->buffers = new RenderBuffers(tile_device);
-				tile->buffers->reset(tile_device, buffer_params);
-				store_rtile = true;
-			}
-			else {
-				assert(rtile.x == stored_rtile.x &&
-				       rtile.y == stored_rtile.y &&
-				       rtile.w == stored_rtile.w &&
-				       rtile.h == stored_rtile.h);
-				tile_lock.unlock();
-				tile->buffers = stored_rtile.buffers;
-			}
-		}
-		else {
-			tile->buffers = new RenderBuffers(tile_device);
-
-			tile->buffers->reset(tile_device, buffer_params);
-		}
+		tile->buffers = new RenderBuffers(tile_device);
+		tile->buffers->reset(tile_device, buffer_params);
 	}
 
 	tile->buffers->params.get_offset_stride(rtile.offset, rtile.stride);
 
 	rtile.buffer = tile->buffers->buffer.device_pointer;
 	rtile.buffers = tile->buffers;
-	rtile.sample = 0;
-
-	if(store_rtile) {
-		render_tiles[tile->index] = rtile;
-		tile_lock.unlock();
-	}
+	rtile.sample = tile_manager.state.sample;
 
 	/* this will tag tile as IN PROGRESS in blender-side render pipeline,
 	 * which is needed to highlight currently rendering tile before first
@@ -484,10 +443,11 @@ void Session::release_tile(RenderTile& rtile)
 	if(tile_manager.finish_tile(rtile.tile_index, delete_tile)) {
 		if(write_render_tile_cb && params.progressive_refine == false) {
 			write_render_tile_cb(rtile);
-			if(delete_tile) {
-				delete rtile.buffers;
-				tile_manager.state.tiles[rtile.tile_index].buffers = NULL;
-			}
+		}
+
+		if(delete_tile) {
+			delete rtile.buffers;
+			tile_manager.state.tiles[rtile.tile_index].buffers = NULL;
 		}
 	}
 	else {
@@ -794,10 +754,10 @@ bool Session::draw(BufferParams& buffer_params, DeviceDrawParams &draw_params)
 
 void Session::reset_(BufferParams& buffer_params, int samples)
 {
-	if(buffers) {
-		if(buffer_params.modified(buffers->params)) {
-			gpu_draw_ready = false;
-			buffers->reset(device, buffer_params);
+	if(buffers && buffer_params.modified(tile_manager.params)) {
+		gpu_draw_ready = false;
+		buffers->reset(device, buffer_params);
+		if(display) {
 			display->reset(device, buffer_params);
 		}
 	}
@@ -819,15 +779,6 @@ void Session::reset(BufferParams& buffer_params, int samples)
 		reset_gpu(buffer_params, samples);
 	else
 		reset_cpu(buffer_params, samples);
-
-	if(params.progressive_refine) {
-		thread_scoped_lock buffers_lock(buffers_mutex);
-
-		foreach(RenderTile &rtile, render_tiles)
-			delete rtile.buffers;
-
-		render_tiles.clear();
-	}
 }
 
 void Session::set_samples(int samples)
@@ -971,7 +922,7 @@ void Session::update_status_time(bool show_pause, bool show_done)
 void Session::render()
 {
 	/* Clear buffers. */
-	if(buffers && tile_manager.state.sample == 0) {
+	if(buffers && tile_manager.state.sample == tile_manager.range_start_sample) {
 		buffers->zero(device);
 	}
 
@@ -1045,8 +996,18 @@ bool Session::update_progressive_refine(bool cancel)
 	}
 
 	if(params.progressive_refine) {
-		foreach(RenderTile &rtile, render_tiles) {
+		foreach(Tile& tile, tile_manager.state.tiles) {
+			if(!tile.buffers) {
+				continue;
+			}
+
+			RenderTile rtile;
+			rtile.x = tile_manager.state.buffer.full_x + tile.x;
+			rtile.y = tile_manager.state.buffer.full_y + tile.y;
+			rtile.w = tile.w;
+			rtile.h = tile.h;
 			rtile.sample = sample;
+			rtile.buffers = tile.buffers;
 
 			if(write) {
 				if(write_render_tile_cb)
@@ -1068,11 +1029,7 @@ void Session::device_free()
 {
 	scene->device_free();
 
-	foreach(RenderTile &tile, render_tiles)
-		delete tile.buffers;
-	tile_manager.free_device();
-
-	render_tiles.clear();
+	tile_manager.device_free();
 
 	/* used from background render only, so no need to
 	 * re-create render/display buffers here
diff --git a/intern/cycles/render/session.h b/intern/cycles/render/session.h
index 980eda0876d..8495d95666b 100644
--- a/intern/cycles/render/session.h
+++ b/intern/cycles/render/session.h
@@ -221,8 +221,6 @@ protected:
 	double last_update_time;
 	bool update_progressive_refine(bool cancel);
 
-	vector<RenderTile> render_tiles;
-
 	DeviceRequestedFeatures get_requested_device_features();
 
 	/* ** Split kernel routines ** */
diff --git a/intern/cycles/render/svm.cpp b/intern/cycles/render/svm.cpp
index 32f89897970..278a8a87b20 100644
--- a/intern/cycles/render/svm.cpp
+++ b/intern/cycles/render/svm.cpp
@@ -48,15 +48,15 @@ void SVMShaderManager::reset(Scene * /*scene*/)
 void SVMShaderManager::device_update_shader(Scene *scene,
                                             Shader *shader,
                                             Progress *progress,
-                                            vector<int4> *global_svm_nodes)
+                                            array<int4> *global_svm_nodes)
 {
 	if(progress->get_cancel()) {
 		return;
 	}
 	assert(shader->graph);
 
-	vector<int4> svm_nodes;
-	svm_nodes.push_back(make_int4(NODE_SHADER_JUMP, 0, 0, 0));
+	array<int4> svm_nodes;
+	svm_nodes.push_back_slow(make_int4(NODE_SHADER_JUMP, 0, 0, 0));
 
 	SVMCompiler::Summary summary;
 	SVMCompiler compiler(scene->shader_manager, scene->image_manager);
@@ -79,12 +79,12 @@ void SVMShaderManager::device_update_shader(Scene *scene,
 	global_svm_nodes->resize(global_nodes_size + svm_nodes.size());
 	
 	/* Offset local SVM nodes to a global address space. */
-	int4& jump_node = global_svm_nodes->at(shader->id);
+	int4& jump_node = (*global_svm_nodes)[shader->id];
 	jump_node.y = svm_nodes[0].y + global_nodes_size - 1;
 	jump_node.z = svm_nodes[0].z + global_nodes_size - 1;
 	jump_node.w = svm_nodes[0].w + global_nodes_size - 1;
 	/* Copy new nodes to global storage. */
-	memcpy(&global_svm_nodes->at(global_nodes_size),
+	memcpy(&(*global_svm_nodes)[global_nodes_size],
 	       &svm_nodes[1],
 	       sizeof(int4) * (svm_nodes.size() - 1));
 	nodes_lock_.unlock();
@@ -106,11 +106,11 @@ void SVMShaderManager::device_update(Device *device, DeviceScene *dscene, Scene
 	device_update_shaders_used(scene);
 
 	/* svm_nodes */
-	vector<int4> svm_nodes;
+	array<int4> svm_nodes;
 	size_t i;
 
 	for(i = 0; i < scene->shaders.size(); i++) {
-		svm_nodes.push_back(make_int4(NODE_SHADER_JUMP, 0, 0, 0));
+		svm_nodes.push_back_slow(make_int4(NODE_SHADER_JUMP, 0, 0, 0));
 	}
 
 	TaskPool task_pool;
@@ -129,7 +129,7 @@ void SVMShaderManager::device_update(Device *device, DeviceScene *dscene, Scene
 		return;
 	}
 
-	dscene->svm_nodes.copy((uint4*)&svm_nodes[0], svm_nodes.size());
+	dscene->svm_nodes.steal_data(svm_nodes);
 	device->tex_alloc("__svm_nodes", dscene->svm_nodes);
 
 	for(i = 0; i < scene->shaders.size(); i++) {
@@ -366,17 +366,17 @@ uint SVMCompiler::encode_uchar4(uint x, uint y, uint z, uint w)
 
 void SVMCompiler::add_node(int a, int b, int c, int d)
 {
-	current_svm_nodes.push_back(make_int4(a, b, c, d));
+	current_svm_nodes.push_back_slow(make_int4(a, b, c, d));
 }
 
 void SVMCompiler::add_node(ShaderNodeType type, int a, int b, int c)
 {
-	current_svm_nodes.push_back(make_int4(type, a, b, c));
+	current_svm_nodes.push_back_slow(make_int4(type, a, b, c));
 }
 
 void SVMCompiler::add_node(ShaderNodeType type, const float3& f)
 {
-	current_svm_nodes.push_back(make_int4(type,
+	current_svm_nodes.push_back_slow(make_int4(type,
 		__float_as_int(f.x),
 		__float_as_int(f.y),
 		__float_as_int(f.z)));
@@ -384,7 +384,7 @@ void SVMCompiler::add_node(ShaderNodeType type, const float3& f)
 
 void SVMCompiler::add_node(const float4& f)
 {
-	current_svm_nodes.push_back(make_int4(
+	current_svm_nodes.push_back_slow(make_int4(
 		__float_as_int(f.x),
 		__float_as_int(f.y),
 		__float_as_int(f.z),
@@ -627,7 +627,7 @@ void SVMCompiler::generate_multi_closure(ShaderNode *root_node,
 				/* Add instruction to skip closure and its dependencies if mix
 				 * weight is zero.
 				 */
-				current_svm_nodes.push_back(make_int4(NODE_JUMP_IF_ONE,
+				current_svm_nodes.push_back_slow(make_int4(NODE_JUMP_IF_ONE,
 				                                      0,
 				                                      stack_assign(facin),
 				                                      0));
@@ -645,7 +645,7 @@ void SVMCompiler::generate_multi_closure(ShaderNode *root_node,
 				/* Add instruction to skip closure and its dependencies if mix
 				 * weight is zero.
 				 */
-				current_svm_nodes.push_back(make_int4(NODE_JUMP_IF_ZERO,
+				current_svm_nodes.push_back_slow(make_int4(NODE_JUMP_IF_ZERO,
 				                                      0,
 				                                      stack_assign(facin),
 				                                      0));
@@ -797,7 +797,7 @@ void SVMCompiler::compile_type(Shader *shader, ShaderGraph *graph, ShaderType ty
 
 void SVMCompiler::compile(Scene *scene,
                           Shader *shader,
-                          vector<int4>& svm_nodes,
+                          array<int4>& svm_nodes,
                           int index,
                           Summary *summary)
 {
@@ -839,9 +839,7 @@ void SVMCompiler::compile(Scene *scene,
 		scoped_timer timer((summary != NULL)? &summary->time_generate_bump: NULL);
 		compile_type(shader, shader->graph, SHADER_TYPE_BUMP);
 		svm_nodes[index].y = svm_nodes.size();
-		svm_nodes.insert(svm_nodes.end(),
-		                 current_svm_nodes.begin(),
-		                 current_svm_nodes.end());
+		svm_nodes.append(current_svm_nodes);
 	}
 
 	/* generate surface shader */
@@ -852,9 +850,7 @@ void SVMCompiler::compile(Scene *scene,
 		if(!has_bump) {
 			svm_nodes[index].y = svm_nodes.size();
 		}
-		svm_nodes.insert(svm_nodes.end(),
-		                 current_svm_nodes.begin(),
-		                 current_svm_nodes.end());
+		svm_nodes.append(current_svm_nodes);
 	}
 
 	/* generate volume shader */
@@ -862,9 +858,7 @@ void SVMCompiler::compile(Scene *scene,
 		scoped_timer timer((summary != NULL)? &summary->time_generate_volume: NULL);
 		compile_type(shader, shader->graph, SHADER_TYPE_VOLUME);
 		svm_nodes[index].z = svm_nodes.size();
-		svm_nodes.insert(svm_nodes.end(),
-		                 current_svm_nodes.begin(),
-		                 current_svm_nodes.end());
+		svm_nodes.append(current_svm_nodes);
 	}
 
 	/* generate displacement shader */
@@ -872,9 +866,7 @@ void SVMCompiler::compile(Scene *scene,
 		scoped_timer timer((summary != NULL)? &summary->time_generate_displacement: NULL);
 		compile_type(shader, shader->graph, SHADER_TYPE_DISPLACEMENT);
 		svm_nodes[index].w = svm_nodes.size();
-		svm_nodes.insert(svm_nodes.end(),
-		                 current_svm_nodes.begin(),
-		                 current_svm_nodes.end());
+		svm_nodes.append(current_svm_nodes);
 	}
 
 	/* Fill in summary information. */
diff --git a/intern/cycles/render/svm.h b/intern/cycles/render/svm.h
index 98ef5fa05d8..0e9905957c7 100644
--- a/intern/cycles/render/svm.h
+++ b/intern/cycles/render/svm.h
@@ -55,7 +55,7 @@ protected:
 	void device_update_shader(Scene *scene,
 	                          Shader *shader,
 	                          Progress *progress,
-	                          vector<int4> *global_svm_nodes);
+	                          array<int4> *global_svm_nodes);
 };
 
 /* Graph Compiler */
@@ -98,7 +98,7 @@ public:
 	SVMCompiler(ShaderManager *shader_manager, ImageManager *image_manager);
 	void compile(Scene *scene,
 	             Shader *shader,
-	             vector<int4>& svm_nodes,
+	             array<int4>& svm_nodes,
 	             int index,
 	             Summary *summary = NULL);
 
@@ -207,7 +207,7 @@ protected:
 	/* compile */
 	void compile_type(Shader *shader, ShaderGraph *graph, ShaderType type);
 
-	vector<int4> current_svm_nodes;
+	array<int4> current_svm_nodes;
 	ShaderType current_type;
 	Shader *current_shader;
 	ShaderGraph *current_graph;
diff --git a/intern/cycles/render/tables.cpp b/intern/cycles/render/tables.cpp
index bf1ef12d602..c08c83cfe11 100644
--- a/intern/cycles/render/tables.cpp
+++ b/intern/cycles/render/tables.cpp
@@ -90,7 +90,9 @@ size_t LookupTables::add_table(DeviceScene *dscene, vector<float>& data)
 	}
 
 	/* copy table data and return offset */
-	dscene->lookup_table.copy_at(&data[0], new_table.offset, data.size());
+	float *dtable = dscene->lookup_table.get_data();
+	memcpy(dtable + new_table.offset, &data[0], sizeof(float) * data.size());
+
 	return new_table.offset;
 }
 
diff --git a/intern/cycles/render/tile.cpp b/intern/cycles/render/tile.cpp
index a9620f79fa0..a388f5dfc8b 100644
--- a/intern/cycles/render/tile.cpp
+++ b/intern/cycles/render/tile.cpp
@@ -17,6 +17,7 @@
 #include "render/tile.h"
 
 #include "util/util_algorithm.h"
+#include "util/util_foreach.h"
 #include "util/util_types.h"
 
 CCL_NAMESPACE_BEGIN
@@ -113,14 +114,16 @@ TileManager::~TileManager()
 {
 }
 
-void TileManager::free_device()
+void TileManager::device_free()
 {
-	if(schedule_denoising) {
+	if(schedule_denoising || progressive) {
 		for(int i = 0; i < state.tiles.size(); i++) {
 			delete state.tiles[i].buffers;
 			state.tiles[i].buffers = NULL;
 		}
 	}
+
+	state.tiles.clear();
 }
 
 static int get_divider(int w, int h, int start_resolution)
@@ -150,7 +153,7 @@ void TileManager::reset(BufferParams& params_, int num_samples_)
 	state.resolution_divider = get_divider(params.width, params.height, start_resolution);
 	state.render_tiles.clear();
 	state.denoising_tiles.clear();
-	state.tiles.clear();
+	device_free();
 }
 
 void TileManager::set_samples(int num_samples_)
@@ -196,7 +199,7 @@ int TileManager::gen_tiles(bool sliced)
 	int slice_num = sliced? num: 1;
 	int tile_w = (tile_size.x >= image_w) ? 1 : divide_up(image_w, tile_size.x);
 
-	state.tiles.clear();
+	device_free();
 	state.render_tiles.clear();
 	state.denoising_tiles.clear();
 	state.render_tiles.resize(num);
@@ -345,6 +348,14 @@ int TileManager::gen_tiles(bool sliced)
 	return idx;
 }
 
+void TileManager::gen_render_tiles()
+{
+	/* Regenerate just the render tiles for progressive render. */
+	foreach(Tile& tile, state.tiles) {
+		state.render_tiles[tile.device].push_back(tile.index);
+	}
+}
+
 void TileManager::set_tiles()
 {
 	int resolution = state.resolution_divider;
@@ -401,6 +412,10 @@ bool TileManager::finish_tile(int index, bool &delete_tile)
 {
 	delete_tile = false;
 
+	if(progressive) {
+		return true;
+	}
+
 	switch(state.tiles[index].state) {
 		case Tile::RENDER:
 		{
@@ -501,7 +516,13 @@ bool TileManager::next()
 			state.num_samples = range_num_samples;
 
 		state.resolution_divider = pixel_size;
-		set_tiles();
+
+		if(state.sample == range_start_sample) {
+			set_tiles();
+		}
+		else {
+			gen_render_tiles();
+		}
 	}
 
 	return true;
diff --git a/intern/cycles/render/tile.h b/intern/cycles/render/tile.h
index 4cd57b7b30c..2692c7cf9f0 100644
--- a/intern/cycles/render/tile.h
+++ b/intern/cycles/render/tile.h
@@ -91,7 +91,7 @@ public:
 	            bool preserve_tile_device, bool background, TileOrder tile_order, int num_devices = 1, int pixel_size = 1);
 	~TileManager();
 
-	void free_device();
+	void device_free();
 	void reset(BufferParams& params, int num_samples);
 	void set_samples(int num_samples);
 	bool next();
@@ -146,6 +146,7 @@ protected:
 
 	/* Generate tile list, return number of tiles. */
 	int gen_tiles(bool sliced);
+	void gen_render_tiles();
 
 	int get_neighbor_index(int index, int neighbor);
 	bool check_neighbor_state(int index, Tile::State state);
diff --git a/intern/cycles/util/util_vector.h b/intern/cycles/util/util_vector.h
index 4add91a3368..9e74505b14a 100644
--- a/intern/cycles/util/util_vector.h
+++ b/intern/cycles/util/util_vector.h
@@ -273,6 +273,15 @@ public:
 		push_back_slow(t);
 	}
 
+	void append(const array<T>& from)
+	{
+		if(from.size()) {
+			size_t old_size = size();
+			resize(old_size + from.size());
+			memcpy(data_ + old_size, from.data(), sizeof(T) * from.size());
+		}
+	}
+
 protected:
 	inline T* mem_allocate(size_t N)
 	{