Code refactor: move more memory allocation logic into device API.

* Remove tex_* and pixels_* functions, replace by mem_*.
* Add MEM_TEXTURE and MEM_PIXELS as memory types recognized by devices.
* No longer create device_memory and call mem_* directly, always go
  through device_only_memory, device_vector and device_pixels.

30 files changed, 934 insertions, 855 deletions

diff --git a/intern/cycles/device/device.cpp b/intern/cycles/device/device.cpp
index 9de10c184fb..41fbe7ce81b 100644
--- a/intern/cycles/device/device.cpp
+++ b/intern/cycles/device/device.cpp
@@ -85,28 +85,12 @@ Device::~Device()
 	}
 }
 
-void Device::pixels_alloc(device_memory& mem)
-{
-	mem_alloc(mem);
-}
-
-void Device::pixels_copy_from(device_memory& mem, int y, int w, int h)
-{
-	if(mem.data_type == TYPE_HALF)
-		mem_copy_from(mem, y, w, h, sizeof(half4));
-	else
-		mem_copy_from(mem, y, w, h, sizeof(uchar4));
-}
-
-void Device::pixels_free(device_memory& mem)
-{
-	mem_free(mem);
-}
-
 void Device::draw_pixels(device_memory& rgba, int y, int w, int h, int dx, int dy, int width, int height, bool transparent,
 	const DeviceDrawParams &draw_params)
 {
-	pixels_copy_from(rgba, y, w, h);
+	assert(mem.type == MEM_PIXELS);
+
+	mem_copy_from(rgba, y, w, h, rgba.memory_elements_size(1));
 
 	if(transparent) {
 		glEnable(GL_BLEND);
diff --git a/intern/cycles/device/device.h b/intern/cycles/device/device.h
index 6bb65cde2a3..316bf70a5c3 100644
--- a/intern/cycles/device/device.h
+++ b/intern/cycles/device/device.h
@@ -281,28 +281,12 @@ public:
 	/* statistics */
 	Stats &stats;
 
-	/* regular memory */
-	virtual void mem_alloc(device_memory& mem) = 0;
-	virtual void mem_copy_to(device_memory& mem) = 0;
-	virtual void mem_copy_from(device_memory& mem,
-		int y, int w, int h, int elem) = 0;
-	virtual void mem_zero(device_memory& mem) = 0;
-	virtual void mem_free(device_memory& mem) = 0;
-
+	/* memory alignment */
 	virtual int mem_address_alignment() { return 16; }
 
 	/* constant memory */
 	virtual void const_copy_to(const char *name, void *host, size_t size) = 0;
 
-	/* texture memory */
-	virtual void tex_alloc(device_memory& /*mem*/) {};
-	virtual void tex_free(device_memory& /*mem*/) {};
-
-	/* pixel memory */
-	virtual void pixels_alloc(device_memory& mem);
-	virtual void pixels_copy_from(device_memory& mem, int y, int w, int h);
-	virtual void pixels_free(device_memory& mem);
-
 	/* open shading language, only for CPU device */
 	virtual void *osl_memory() { return NULL; }
 
@@ -349,6 +333,20 @@ public:
 	static void tag_update();
 
 	static void free_memory();
+
+protected:
+	/* Memory allocation, only accessed through device_memory. */
+	friend class MultiDevice;
+	friend class DeviceServer;
+	friend class device_memory;
+
+	virtual void mem_alloc(device_memory& mem) = 0;
+	virtual void mem_copy_to(device_memory& mem) = 0;
+	virtual void mem_copy_from(device_memory& mem,
+		int y, int w, int h, int elem) = 0;
+	virtual void mem_zero(device_memory& mem) = 0;
+	virtual void mem_free(device_memory& mem) = 0;
+
 private:
 	/* Indicted whether device types and devices lists were initialized. */
 	static bool need_types_update, need_devices_update;
diff --git a/intern/cycles/device/device_cpu.cpp b/intern/cycles/device/device_cpu.cpp
index b4398f21014..32ab18fe164 100644
--- a/intern/cycles/device/device_cpu.cpp
+++ b/intern/cycles/device/device_cpu.cpp
@@ -209,7 +209,7 @@ public:
 
 	CPUDevice(DeviceInfo& info_, Stats &stats_, bool background_)
 	: Device(info_, stats_, background_),
-	  texture_info(this, "__texture_info"),
+	  texture_info(this, "__texture_info", MEM_TEXTURE),
 #define REGISTER_KERNEL(name) name ## _kernel(KERNEL_FUNCTIONS(name))
 	  REGISTER_KERNEL(path_trace),
 	  REGISTER_KERNEL(convert_to_half_float),
@@ -269,7 +269,7 @@ public:
 	~CPUDevice()
 	{
 		task_pool.stop();
-		tex_free(texture_info);
+		texture_info.free();
 	}
 
 	virtual bool show_samples() const
@@ -280,33 +280,50 @@ public:
 	void load_texture_info()
 	{
 		if(need_texture_info) {
-			tex_free(texture_info);
-			tex_alloc(texture_info);
+			texture_info.copy_to_device();
 			need_texture_info = false;
 		}
 	}
 
 	void mem_alloc(device_memory& mem)
 	{
-		if(mem.name) {
-			VLOG(1) << "Buffer allocate: " << mem.name << ", "
-			        << string_human_readable_number(mem.memory_size()) << " bytes. ("
-			        << string_human_readable_size(mem.memory_size()) << ")";
+		if(mem.type == MEM_TEXTURE) {
+			assert(!"mem_alloc not supported for textures.");
 		}
+		else {
+			if(mem.name) {
+				VLOG(1) << "Buffer allocate: " << mem.name << ", "
+						<< string_human_readable_number(mem.memory_size()) << " bytes. ("
+						<< string_human_readable_size(mem.memory_size()) << ")";
+			}
 
-		mem.device_pointer = mem.data_pointer;
+			mem.device_pointer = mem.data_pointer;
 
-		if(!mem.device_pointer) {
-			mem.device_pointer = (device_ptr)malloc(mem.memory_size());
-		}
+			if(!mem.device_pointer) {
+				mem.device_pointer = (device_ptr)malloc(mem.memory_size());
+			}
 
-		mem.device_size = mem.memory_size();
-		stats.mem_alloc(mem.device_size);
+			mem.device_size = mem.memory_size();
+			stats.mem_alloc(mem.device_size);
+		}
 	}
 
-	void mem_copy_to(device_memory& /*mem*/)
+	void mem_copy_to(device_memory& mem)
 	{
-		/* no-op */
+		if(mem.type == MEM_TEXTURE) {
+			tex_free(mem);
+			tex_alloc(mem);
+		}
+		else if(mem.type == MEM_PIXELS) {
+			assert(!"mem_copy_to not supported for pixels.");
+		}
+		else {
+			if(!mem.device_pointer) {
+				mem_alloc(mem);
+			}
+
+			/* copy is no-op */
+		}
 	}
 
 	void mem_copy_from(device_memory& /*mem*/,
@@ -318,12 +335,21 @@ public:
 
 	void mem_zero(device_memory& mem)
 	{
-		memset((void*)mem.device_pointer, 0, mem.memory_size());
+		if(!mem.device_pointer) {
+			mem_alloc(mem);
+		}
+
+		if(mem.device_pointer) {
+			memset((void*)mem.device_pointer, 0, mem.memory_size());
+		}
 	}
 
 	void mem_free(device_memory& mem)
 	{
-		if(mem.device_pointer) {
+		if(mem.type == MEM_TEXTURE) {
+			tex_free(mem);
+		}
+		else if(mem.device_pointer) {
 			if(!mem.data_pointer) {
 				free((void*)mem.device_pointer);
 			}
@@ -354,7 +380,7 @@ public:
 			kernel_tex_copy(&kernel_globals,
 							mem.name,
 							mem.data_pointer,
-							mem.data_width);
+							mem.data_size);
 		}
 		else {
 			/* Image Texture. */
@@ -431,13 +457,13 @@ public:
 
 	bool denoising_set_tiles(device_ptr *buffers, DenoisingTask *task)
 	{
-		mem_alloc(task->tiles_mem);
-
 		TilesInfo *tiles = (TilesInfo*) task->tiles_mem.data_pointer;
 		for(int i = 0; i < 9; i++) {
 			tiles->buffers[i] = buffers[i];
 		}
 
+		task->tiles_mem.copy_to_device();
+
 		return true;
 	}
 
@@ -723,8 +749,7 @@ public:
 
 		/* allocate buffer for kernel globals */
 		device_only_memory<KernelGlobals> kgbuffer(this, "kernel_globals");
-		kgbuffer.resize(1);
-		mem_alloc(kgbuffer);
+		kgbuffer.alloc_to_device(1);
 
 		KernelGlobals *kg = new ((void*) kgbuffer.device_pointer) KernelGlobals(thread_kernel_globals_init());
 
@@ -734,8 +759,7 @@ public:
 			requested_features.max_closure = MAX_CLOSURE;
 			if(!split_kernel->load_kernels(requested_features)) {
 				thread_kernel_globals_free((KernelGlobals*)kgbuffer.device_pointer);
-				mem_free(kgbuffer);
-
+				kgbuffer.free();
 				delete split_kernel;
 				return;
 			}
@@ -766,7 +790,7 @@ public:
 
 		thread_kernel_globals_free((KernelGlobals*)kgbuffer.device_pointer);
 		kg->~KernelGlobals();
-		mem_free(kgbuffer);
+		kgbuffer.free();
 		delete split_kernel;
 	}
 
diff --git a/intern/cycles/device/device_cuda.cpp b/intern/cycles/device/device_cuda.cpp
index be606a92434..aa6386e455b 100644
--- a/intern/cycles/device/device_cuda.cpp
+++ b/intern/cycles/device/device_cuda.cpp
@@ -218,7 +218,7 @@ public:
 
 	CUDADevice(DeviceInfo& info, Stats &stats, bool background_)
 	: Device(info, stats, background_),
-	  texture_info(this, "__texture_info")
+	  texture_info(this, "__texture_info", MEM_TEXTURE)
 	{
 		first_error = true;
 		background = background_;
@@ -275,7 +275,7 @@ public:
 		delete split_kernel;
 
 		if(info.has_bindless_textures) {
-			tex_free(texture_info);
+			texture_info.free();
 		}
 
 		cuda_assert(cuCtxDestroy(cuContext));
@@ -548,20 +548,19 @@ public:
 	void load_texture_info()
 	{
 		if(info.has_bindless_textures && need_texture_info) {
-			tex_free(texture_info);
-			tex_alloc(texture_info);
+			texture_info.copy_to_device();
 			need_texture_info = false;
 		}
 	}
 
-	void mem_alloc(device_memory& mem)
+	void generic_alloc(device_memory& mem)
 	{
 		CUDAContextScope scope(this);
 
 		if(mem.name) {
 			VLOG(1) << "Buffer allocate: " << mem.name << ", "
-			        << string_human_readable_number(mem.memory_size()) << " bytes. ("
-			        << string_human_readable_size(mem.memory_size()) << ")";
+					<< string_human_readable_number(mem.memory_size()) << " bytes. ("
+					<< string_human_readable_size(mem.memory_size()) << ")";
 		}
 
 		CUdeviceptr device_pointer;
@@ -572,31 +571,88 @@ public:
 		stats.mem_alloc(size);
 	}
 
+	void generic_copy_to(device_memory& mem)
+	{
+		if(mem.device_pointer) {
+			CUDAContextScope scope(this);
+			cuda_assert(cuMemcpyHtoD(cuda_device_ptr(mem.device_pointer), (void*)mem.data_pointer, mem.memory_size()));
+		}
+	}
+
+	void generic_free(device_memory& mem)
+	{
+		if(mem.device_pointer) {
+			CUDAContextScope scope(this);
+
+			cuda_assert(cuMemFree(cuda_device_ptr(mem.device_pointer)));
+
+			mem.device_pointer = 0;
+
+			stats.mem_free(mem.device_size);
+			mem.device_size = 0;
+		}
+	}
+
+	void mem_alloc(device_memory& mem)
+	{
+		if(mem.type == MEM_PIXELS && !background) {
+			pixels_alloc(mem);
+		}
+		else if(mem.type == MEM_TEXTURE) {
+			assert(!"mem_alloc not supported for textures.");
+		}
+		else {
+			generic_alloc(mem);
+		}
+	}
+
 	void mem_copy_to(device_memory& mem)
 	{
-		CUDAContextScope scope(this);
+		if(mem.type == MEM_PIXELS) {
+			assert(!"mem_copy_to not supported for pixels.");
+		}
+		else if(mem.type == MEM_TEXTURE) {
+			tex_free(mem);
+			tex_alloc(mem);
+		}
+		else {
+			if(!mem.device_pointer) {
+				generic_alloc(mem);
+			}
 
-		if(mem.device_pointer)
-			cuda_assert(cuMemcpyHtoD(cuda_device_ptr(mem.device_pointer), (void*)mem.data_pointer, mem.memory_size()));
+			generic_copy_to(mem);
+		}
 	}
 
 	void mem_copy_from(device_memory& mem, int y, int w, int h, int elem)
 	{
-		CUDAContextScope scope(this);
-		size_t offset = elem*y*w;
-		size_t size = elem*w*h;
-
-		if(mem.device_pointer) {
-			cuda_assert(cuMemcpyDtoH((uchar*)mem.data_pointer + offset,
-			                         (CUdeviceptr)(mem.device_pointer + offset), size));
+		if(mem.type == MEM_PIXELS && !background) {
+			pixels_copy_from(mem, y, w, h);
+		}
+		else if(mem.type == MEM_TEXTURE) {
+			assert(!"mem_copy_from not supported for textures.");
 		}
 		else {
-			memset((char*)mem.data_pointer + offset, 0, size);
+			CUDAContextScope scope(this);
+			size_t offset = elem*y*w;
+			size_t size = elem*w*h;
+
+			if(mem.device_pointer) {
+				cuda_assert(cuMemcpyDtoH((uchar*)mem.data_pointer + offset,
+										 (CUdeviceptr)(mem.device_pointer + offset), size));
+			}
+			else {
+				memset((char*)mem.data_pointer + offset, 0, size);
+			}
 		}
 	}
 
 	void mem_zero(device_memory& mem)
 	{
+		if(!mem.device_pointer) {
+			mem_alloc(mem);
+		}
+
 		if(mem.data_pointer) {
 			memset((void*)mem.data_pointer, 0, mem.memory_size());
 		}
@@ -609,14 +665,14 @@ public:
 
 	void mem_free(device_memory& mem)
 	{
-		if(mem.device_pointer) {
-			CUDAContextScope scope(this);
-			cuda_assert(cuMemFree(cuda_device_ptr(mem.device_pointer)));
-
-			mem.device_pointer = 0;
-
-			stats.mem_free(mem.device_size);
-			mem.device_size = 0;
+		if(mem.type == MEM_PIXELS && !background) {
+			pixels_free(mem);
+		}
+		else if(mem.type == MEM_TEXTURE) {
+			tex_free(mem);
+		}
+		else {
+			generic_free(mem);
 		}
 	}
 
@@ -700,8 +756,8 @@ public:
 
 		if(mem.interpolation == INTERPOLATION_NONE) {
 			/* Data Storage */
-			mem_alloc(mem);
-			mem_copy_to(mem);
+			generic_alloc(mem);
+			generic_copy_to(mem);
 
 			CUdeviceptr cumem;
 			size_t cubytes;
@@ -891,21 +947,19 @@ public:
 			}
 			else {
 				tex_interp_map.erase(tex_interp_map.find(mem.device_pointer));
-				mem_free(mem);
+				generic_free(mem);
 			}
 		}
 	}
 
 	bool denoising_set_tiles(device_ptr *buffers, DenoisingTask *task)
 	{
-		mem_alloc(task->tiles_mem);
-
 		TilesInfo *tiles = (TilesInfo*) task->tiles_mem.data_pointer;
 		for(int i = 0; i < 9; i++) {
 			tiles->buffers[i] = buffers[i];
 		}
 
-		mem_copy_to(task->tiles_mem);
+		task->tiles_mem.copy_to_device();
 
 		return !have_error();
 	}
@@ -1272,7 +1326,7 @@ public:
 		task.unmap_neighbor_tiles(rtiles, this);
 	}
 
-	void path_trace(DeviceTask& task, RenderTile& rtile)
+	void path_trace(DeviceTask& task, RenderTile& rtile, device_vector<WorkTile>& work_tiles)
 	{
 		if(have_error())
 			return;
@@ -1295,8 +1349,7 @@ public:
 		cuda_assert(cuFuncSetCacheConfig(cuPathTrace, CU_FUNC_CACHE_PREFER_L1));
 
 		/* Allocate work tile. */
-		device_vector<WorkTile> work_tiles(this, "work_tiles", MEM_READ_ONLY);
-		work_tiles.resize(1);
+		work_tiles.alloc(1);
 
 		WorkTile *wtile = work_tiles.get_data();
 		wtile->x = rtile.x;
@@ -1306,9 +1359,6 @@ public:
 		wtile->offset = rtile.offset;
 		wtile->stride = rtile.stride;
 		wtile->buffer = (float*)cuda_device_ptr(rtile.buffer);
-		mem_alloc(work_tiles);
-
-		CUdeviceptr d_work_tiles = cuda_device_ptr(work_tiles.device_pointer);
 
 		/* Prepare work size. More step samples render faster, but for now we
 		 * remain conservative for GPUs connected to a display to avoid driver
@@ -1329,8 +1379,9 @@ public:
 			/* Setup and copy work tile to device. */
 			wtile->start_sample = sample;
 			wtile->num_samples = min(step_samples, end_sample - sample);;
-			mem_copy_to(work_tiles);
+			work_tiles.copy_to_device();
 
+			CUdeviceptr d_work_tiles = cuda_device_ptr(work_tiles.device_pointer);
 			uint total_work_size = wtile->w * wtile->h * wtile->num_samples;
 			uint num_blocks = divide_up(total_work_size, num_threads_per_block);
 
@@ -1354,8 +1405,6 @@ public:
 					break;
 			}
 		}
-
-		mem_free(work_tiles);
 	}
 
 	void film_convert(DeviceTask& task, device_ptr buffer, device_ptr rgba_byte, device_ptr rgba_half)
@@ -1508,104 +1557,90 @@ public:
 
 	void pixels_alloc(device_memory& mem)
 	{
-		if(!background) {
-			PixelMem pmem;
-
-			pmem.w = mem.data_width;
-			pmem.h = mem.data_height;
+		PixelMem pmem;
 
-			CUDAContextScope scope(this);
+		pmem.w = mem.data_width;
+		pmem.h = mem.data_height;
 
-			glGenBuffers(1, &pmem.cuPBO);
-			glBindBuffer(GL_PIXEL_UNPACK_BUFFER, pmem.cuPBO);
-			if(mem.data_type == TYPE_HALF)
-				glBufferData(GL_PIXEL_UNPACK_BUFFER, pmem.w*pmem.h*sizeof(GLhalf)*4, NULL, GL_DYNAMIC_DRAW);
-			else
-				glBufferData(GL_PIXEL_UNPACK_BUFFER, pmem.w*pmem.h*sizeof(uint8_t)*4, NULL, GL_DYNAMIC_DRAW);
+		CUDAContextScope scope(this);
 
-			glBindBuffer(GL_PIXEL_UNPACK_BUFFER, 0);
+		glGenBuffers(1, &pmem.cuPBO);
+		glBindBuffer(GL_PIXEL_UNPACK_BUFFER, pmem.cuPBO);
+		if(mem.data_type == TYPE_HALF)
+			glBufferData(GL_PIXEL_UNPACK_BUFFER, pmem.w*pmem.h*sizeof(GLhalf)*4, NULL, GL_DYNAMIC_DRAW);
+		else
+			glBufferData(GL_PIXEL_UNPACK_BUFFER, pmem.w*pmem.h*sizeof(uint8_t)*4, NULL, GL_DYNAMIC_DRAW);
 
-			glGenTextures(1, &pmem.cuTexId);
-			glBindTexture(GL_TEXTURE_2D, pmem.cuTexId);
-			if(mem.data_type == TYPE_HALF)
-				glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA16F_ARB, pmem.w, pmem.h, 0, GL_RGBA, GL_HALF_FLOAT, NULL);
-			else
-				glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA8, pmem.w, pmem.h, 0, GL_RGBA, GL_UNSIGNED_BYTE, NULL);
-			glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_NEAREST);
-			glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, GL_NEAREST);
-			glBindTexture(GL_TEXTURE_2D, 0);
+		glBindBuffer(GL_PIXEL_UNPACK_BUFFER, 0);
 
-			CUresult result = cuGraphicsGLRegisterBuffer(&pmem.cuPBOresource, pmem.cuPBO, CU_GRAPHICS_MAP_RESOURCE_FLAGS_NONE);
+		glGenTextures(1, &pmem.cuTexId);
+		glBindTexture(GL_TEXTURE_2D, pmem.cuTexId);
+		if(mem.data_type == TYPE_HALF)
+			glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA16F_ARB, pmem.w, pmem.h, 0, GL_RGBA, GL_HALF_FLOAT, NULL);
+		else
+			glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA8, pmem.w, pmem.h, 0, GL_RGBA, GL_UNSIGNED_BYTE, NULL);
+		glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_NEAREST);
+		glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, GL_NEAREST);
+		glBindTexture(GL_TEXTURE_2D, 0);
 
-			if(result == CUDA_SUCCESS) {
-				mem.device_pointer = pmem.cuTexId;
-				pixel_mem_map[mem.device_pointer] = pmem;
+		CUresult result = cuGraphicsGLRegisterBuffer(&pmem.cuPBOresource, pmem.cuPBO, CU_GRAPHICS_MAP_RESOURCE_FLAGS_NONE);
 
-				mem.device_size = mem.memory_size();
-				stats.mem_alloc(mem.device_size);
+		if(result == CUDA_SUCCESS) {
+			mem.device_pointer = pmem.cuTexId;
+			pixel_mem_map[mem.device_pointer] = pmem;
 
-				return;
-			}
-			else {
-				/* failed to register buffer, fallback to no interop */
-				glDeleteBuffers(1, &pmem.cuPBO);
-				glDeleteTextures(1, &pmem.cuTexId);
+			mem.device_size = mem.memory_size();
+			stats.mem_alloc(mem.device_size);
 
-				background = true;
-			}
+			return;
 		}
+		else {
+			/* failed to register buffer, fallback to no interop */
+			glDeleteBuffers(1, &pmem.cuPBO);
+			glDeleteTextures(1, &pmem.cuTexId);
 
-		Device::pixels_alloc(mem);
+			background = true;
+		}
 	}
 
 	void pixels_copy_from(device_memory& mem, int y, int w, int h)
 	{
-		if(!background) {
-			PixelMem pmem = pixel_mem_map[mem.device_pointer];
-
-			CUDAContextScope scope(this);
-
-			glBindBuffer(GL_PIXEL_UNPACK_BUFFER, pmem.cuPBO);
-			uchar *pixels = (uchar*)glMapBuffer(GL_PIXEL_UNPACK_BUFFER, GL_READ_ONLY);
-			size_t offset = sizeof(uchar)*4*y*w;
-			memcpy((uchar*)mem.data_pointer + offset, pixels + offset, sizeof(uchar)*4*w*h);
-			glUnmapBuffer(GL_PIXEL_UNPACK_BUFFER);
-			glBindBuffer(GL_PIXEL_UNPACK_BUFFER, 0);
+		PixelMem pmem = pixel_mem_map[mem.device_pointer];
 
-			return;
-		}
+		CUDAContextScope scope(this);
 
-		Device::pixels_copy_from(mem, y, w, h);
+		glBindBuffer(GL_PIXEL_UNPACK_BUFFER, pmem.cuPBO);
+		uchar *pixels = (uchar*)glMapBuffer(GL_PIXEL_UNPACK_BUFFER, GL_READ_ONLY);
+		size_t offset = sizeof(uchar)*4*y*w;
+		memcpy((uchar*)mem.data_pointer + offset, pixels + offset, sizeof(uchar)*4*w*h);
+		glUnmapBuffer(GL_PIXEL_UNPACK_BUFFER);
+		glBindBuffer(GL_PIXEL_UNPACK_BUFFER, 0);
 	}
 
 	void pixels_free(device_memory& mem)
 	{
 		if(mem.device_pointer) {
-			if(!background) {
-				PixelMem pmem = pixel_mem_map[mem.device_pointer];
-
-				CUDAContextScope scope(this);
-
-				cuda_assert(cuGraphicsUnregisterResource(pmem.cuPBOresource));
-				glDeleteBuffers(1, &pmem.cuPBO);
-				glDeleteTextures(1, &pmem.cuTexId);
+			PixelMem pmem = pixel_mem_map[mem.device_pointer];
 
-				pixel_mem_map.erase(pixel_mem_map.find(mem.device_pointer));
-				mem.device_pointer = 0;
+			CUDAContextScope scope(this);
 
-				stats.mem_free(mem.device_size);
-				mem.device_size = 0;
+			cuda_assert(cuGraphicsUnregisterResource(pmem.cuPBOresource));
+			glDeleteBuffers(1, &pmem.cuPBO);
+			glDeleteTextures(1, &pmem.cuTexId);
 
-				return;
-			}
+			pixel_mem_map.erase(pixel_mem_map.find(mem.device_pointer));
+			mem.device_pointer = 0;
 
-			Device::pixels_free(mem);
+			stats.mem_free(mem.device_size);
+			mem.device_size = 0;
 		}
 	}
 
 	void draw_pixels(device_memory& mem, int y, int w, int h, int dx, int dy, int width, int height, bool transparent,
 		const DeviceDrawParams &draw_params)
 	{
+		assert(mem.type == MEM_PIXELS);
+
 		if(!background) {
 			PixelMem pmem = pixel_mem_map[mem.device_pointer];
 			float *vpointer;
@@ -1724,6 +1759,8 @@ public:
 				}
 			}
 
+			device_vector<WorkTile> work_tiles(this, "work_tiles", MEM_READ_ONLY);
+
 			/* keep rendering tiles until done */
 			while(task->acquire_tile(this, tile)) {
 				if(tile.task == RenderTile::PATH_TRACE) {
@@ -1732,7 +1769,7 @@ public:
 						split_kernel->path_trace(task, tile, void_buffer, void_buffer);
 					}
 					else {
-						path_trace(*task, tile);
+						path_trace(*task, tile, work_tiles);
 					}
 				}
 				else if(tile.task == RenderTile::DENOISE) {
@@ -1750,6 +1787,8 @@ public:
 						break;
 				}
 			}
+
+			work_tiles.free();
 		}
 		else if(task->type == DeviceTask::SHADER) {
 			shader(*task);
@@ -1884,8 +1923,8 @@ uint64_t CUDASplitKernel::state_buffer_size(device_memory& /*kg*/, device_memory
 	CUDAContextScope scope(device);
 
 	device_vector<uint64_t> size_buffer(device, "size_buffer", MEM_READ_WRITE);
-	size_buffer.resize(1);
-	device->mem_alloc(size_buffer);
+	size_buffer.alloc(1);
+	size_buffer.zero_to_device();
 
 	uint threads = num_threads;
 	CUdeviceptr d_size = device->cuda_device_ptr(size_buffer.device_pointer);
@@ -1908,9 +1947,9 @@ uint64_t CUDASplitKernel::state_buffer_size(device_memory& /*kg*/, device_memory
 	                           1, 1, 1,
 	                           0, 0, (void**)&args, 0));
 
-	device->mem_copy_from(size_buffer, 0, 1, 1, sizeof(uint64_t));
+	size_buffer.copy_from_device(0, 1, 1);
 	size_t size = size_buffer[0];
-	device->mem_free(size_buffer);
+	size_buffer.free();
 
 	return size;
 }
diff --git a/intern/cycles/device/device_denoising.cpp b/intern/cycles/device/device_denoising.cpp
index 2c3bfefd8b0..2d39721e3d3 100644
--- a/intern/cycles/device/device_denoising.cpp
+++ b/intern/cycles/device/device_denoising.cpp
@@ -44,7 +44,7 @@ void DenoisingTask::init_from_devicetask(const DeviceTask &task)
 
 void DenoisingTask::tiles_from_rendertiles(RenderTile *rtiles)
 {
-	tiles = (TilesInfo*) tiles_mem.resize(sizeof(TilesInfo)/sizeof(int));
+	tiles = (TilesInfo*) tiles_mem.alloc(sizeof(TilesInfo)/sizeof(int));
 
 	device_ptr buffers[9];
 	for(int i = 0; i < 9; i++) {
@@ -75,8 +75,7 @@ bool DenoisingTask::run_denoising()
 	buffer.w = align_up(rect.z - rect.x, 4);
 	buffer.h = rect.w - rect.y;
 	buffer.pass_stride = align_up(buffer.w * buffer.h, divide_up(device->mem_address_alignment(), sizeof(float)));
-	buffer.mem.resize(buffer.pass_stride * buffer.passes);
-	device->mem_alloc(buffer.mem);
+	buffer.mem.alloc_to_device(buffer.pass_stride * buffer.passes);
 
 	device_ptr null_ptr = (device_ptr) 0;
 
@@ -161,8 +160,7 @@ bool DenoisingTask::run_denoising()
 		int num_color_passes = 3;
 
 		device_only_memory<float> temp_color(device, "Denoising temporary color");
-		temp_color.resize(3*buffer.pass_stride);
-		device->mem_alloc(temp_color);
+		temp_color.alloc_to_device(3*buffer.pass_stride);
 
 		for(int pass = 0; pass < num_color_passes; pass++) {
 			device_sub_ptr color_pass(temp_color, pass*buffer.pass_stride, buffer.pass_stride);
@@ -177,31 +175,25 @@ bool DenoisingTask::run_denoising()
 			functions.detect_outliers(temp_color.device_pointer, *color_var_pass, *depth_pass, *output_pass);
 		}
 
-		device->mem_free(temp_color);
+		temp_color.free();
 	}
 
 	storage.w = filter_area.z;
 	storage.h = filter_area.w;
-	storage.transform.resize(storage.w*storage.h*TRANSFORM_SIZE);
-	storage.rank.resize(storage.w*storage.h);
-	device->mem_alloc(storage.transform);
-	device->mem_alloc(storage.rank);
+	storage.transform.alloc_to_device(storage.w*storage.h*TRANSFORM_SIZE);
+	storage.rank.alloc_to_device(storage.w*storage.h);
 
 	functions.construct_transform();
 
 	device_only_memory<float> temporary_1(device, "Denoising NLM temporary 1");
 	device_only_memory<float> temporary_2(device, "Denoising NLM temporary 2");
-	temporary_1.resize(buffer.w*buffer.h);
-	temporary_2.resize(buffer.w*buffer.h);
-	device->mem_alloc(temporary_1);
-	device->mem_alloc(temporary_2);
+	temporary_1.alloc_to_device(buffer.w*buffer.h);
+	temporary_2.alloc_to_device(buffer.w*buffer.h);
 	reconstruction_state.temporary_1_ptr = temporary_1.device_pointer;
 	reconstruction_state.temporary_2_ptr = temporary_2.device_pointer;
 
-	storage.XtWX.resize(storage.w*storage.h*XTWX_SIZE);
-	storage.XtWY.resize(storage.w*storage.h*XTWY_SIZE);
-	device->mem_alloc(storage.XtWX);
-	device->mem_alloc(storage.XtWY);
+	storage.XtWX.alloc_to_device(storage.w*storage.h*XTWX_SIZE);
+	storage.XtWY.alloc_to_device(storage.w*storage.h*XTWY_SIZE);
 
 	reconstruction_state.filter_rect = make_int4(filter_area.x-rect.x, filter_area.y-rect.y, storage.w, storage.h);
 	int tile_coordinate_offset = filter_area.y*render_buffer.stride + filter_area.x;
@@ -218,14 +210,14 @@ bool DenoisingTask::run_denoising()
 		functions.reconstruct(*color_ptr, *color_var_ptr, render_buffer.ptr);
 	}
 
-	device->mem_free(storage.XtWX);
-	device->mem_free(storage.XtWY);
-	device->mem_free(storage.transform);
-	device->mem_free(storage.rank);
-	device->mem_free(temporary_1);
-	device->mem_free(temporary_2);
-	device->mem_free(buffer.mem);
-	device->mem_free(tiles_mem);
+	storage.XtWX.free();
+	storage.XtWY.free();
+	storage.transform.free();
+	storage.rank.free();
+	temporary_1.free();
+	temporary_2.free();
+	buffer.mem.free();
+	tiles_mem.free();
 	return true;
 }
 
diff --git a/intern/cycles/device/device_memory.cpp b/intern/cycles/device/device_memory.cpp
index 98fa638ef8e..9f4f60e7531 100644
--- a/intern/cycles/device/device_memory.cpp
+++ b/intern/cycles/device/device_memory.cpp
@@ -43,6 +43,68 @@ device_memory::~device_memory()
 {
 }
 
+device_ptr device_memory::host_alloc(size_t size)
+{
+	if(!size) {
+		return 0;
+	}
+
+	size_t alignment = device->mem_address_alignment();
+	device_ptr ptr = (device_ptr)util_aligned_malloc(size, alignment);
+
+	if(ptr) {
+		util_guarded_mem_alloc(size);
+	}
+	else {
+		throw std::bad_alloc();
+	}
+
+	return ptr;
+}
+
+void device_memory::host_free(device_ptr ptr, size_t size)
+{
+	if(ptr) {
+		util_guarded_mem_free(size);
+		util_aligned_free((void*)ptr);
+	}
+}
+
+void device_memory::device_alloc()
+{
+	assert(!device_pointer && type != MEM_TEXTURE);
+	device->mem_alloc(*this);
+}
+
+void device_memory::device_free()
+{
+	if(device_pointer) {
+		device->mem_free(*this);
+	}
+}
+
+void device_memory::device_copy_to()
+{
+	assert(type != MEM_PIXELS && type != MEM_WRITE_ONLY);
+	if(data_size) {
+		device->mem_copy_to(*this);
+	}
+}
+
+void device_memory::device_copy_from(int y, int w, int h, int elem)
+{
+	assert(type != MEM_TEXTURE && type != MEM_READ_ONLY);
+	device->mem_copy_from(*this, y, w, h, elem);
+}
+
+void device_memory::device_zero()
+{
+	assert(type != MEM_PIXELS && type != MEM_WRITE_ONLY);
+	if(data_size) {
+		device->mem_zero(*this);
+	}
+}
+
 /* Device Sub Ptr */
 
 device_sub_ptr::device_sub_ptr(device_memory& mem, int offset, int size)
diff --git a/intern/cycles/device/device_memory.h b/intern/cycles/device/device_memory.h
index 3dfecde59d8..7bf8bdc1cea 100644
--- a/intern/cycles/device/device_memory.h
+++ b/intern/cycles/device/device_memory.h
@@ -19,14 +19,7 @@
 
 /* Device Memory
  *
- * This file defines data types that can be used in device memory arrays, and
- * a device_vector<T> type to store such arrays.
- *
- * device_vector<T> contains an STL vector, metadata about the data type,
- * dimensions, elements, and a device pointer. For the CPU device this is just
- * a pointer to the STL vector data, as no copying needs to take place. For
- * other devices this is a pointer to device memory, where we will copy memory
- * to and from. */
+ * Data types for allocating, copying and freeing device memory. */
 
 #include "util/util_debug.h"
 #include "util/util_half.h"
@@ -41,7 +34,9 @@ class Device;
 enum MemoryType {
 	MEM_READ_ONLY,
 	MEM_WRITE_ONLY,
-	MEM_READ_WRITE
+	MEM_READ_WRITE,
+	MEM_TEXTURE,
+	MEM_PIXELS
 };
 
 /* Supported Data Types */
@@ -172,7 +167,10 @@ template<> struct device_type_traits<uint64_t> {
 	static const int num_elements = 1;
 };
 
-/* Device Memory */
+/* Device Memory
+ *
+ * Base class for all device memory. This should not be allocated directly,
+ * instead the appropriate subclass can be used. */
 
 class device_memory
 {
@@ -182,7 +180,7 @@ public:
 		return elements*data_elements*datatype_size(data_type);
 	}
 
-	/* data information */
+	/* Data information. */
 	DataType data_type;
 	int data_elements;
 	device_ptr data_pointer;
@@ -196,25 +194,39 @@ public:
 	InterpolationType interpolation;
 	ExtensionType extension;
 
-	/* device pointer */
+	/* Device pointer. */
 	Device *device;
 	device_ptr device_pointer;
 
-	device_memory(Device *device, const char *name, MemoryType type);
 	virtual ~device_memory();
 
-	void resize(size_t size)
-	{
-		data_size = size;
-		data_width = size;
-	}
-
 protected:
-	/* no copying */
+	/* Only create through subclasses. */
+	device_memory(Device *device, const char *name, MemoryType type);
+
+	/* No copying allowed. */
 	device_memory(const device_memory&);
 	device_memory& operator = (const device_memory&);
+
+	/* Host allocation on the device. All data_pointer memory should be
+	 * allocated with these functions, for devices that support using
+	 * the same pointer for host and device. */
+	device_ptr host_alloc(size_t size);
+	void host_free(device_ptr ptr, size_t size);
+
+	/* Device memory allocation and copying. */
+	void device_alloc();
+	void device_free();
+	void device_copy_to();
+	void device_copy_from(int y, int w, int h, int elem);
+	void device_zero();
 };
 
+/* Device Only Memory
+ *
+ * Working memory only needed by the device, with no corresponding allocation
+ * on the host. Only used internally in the device implementations. */
+
 template<typename T>
 class device_only_memory : public device_memory
 {
@@ -226,18 +238,43 @@ public:
 		data_elements = max(device_type_traits<T>::num_elements, 1);
 	}
 
-	void resize(size_t num)
+	virtual ~device_only_memory()
+	{
+		free();
+	}
+
+	void alloc_to_device(size_t num)
+	{
+		data_size = num*sizeof(T);
+		device_alloc();
+	}
+
+	void free()
+	{
+		device_free();
+	}
+
+	void zero_to_device()
 	{
-		device_memory::resize(num*sizeof(T));
+		device_zero();
 	}
 };
 
-/* Device Vector */
+/* Device Vector
+ *
+ * Data vector to exchange data between host and device. Memory will be
+ * allocated on the host first with alloc() and resize, and then filled
+ * in and copied to the device with copy_to_device(). Or alternatively
+ * allocated and set to zero on the device with zero_to_device().
+ *
+ * When using memory type MEM_TEXTURE, a pointer to this memory will be
+ * automatically attached to kernel globals, using the provided name
+ * matching an entry in kernel_textures.h. */
 
 template<typename T> class device_vector : public device_memory
 {
 public:
-	device_vector(Device *device, const char *name, MemoryType type = MEM_READ_ONLY)
+	device_vector(Device *device, const char *name, MemoryType type)
 	: device_memory(device, name, type)
 	{
 		data_type = device_type_traits<T>::data_type;
@@ -246,84 +283,175 @@ public:
 		assert(data_elements > 0);
 	}
 
-	virtual ~device_vector() {}
+	virtual ~device_vector()
+	{
+		free();
+	}
 
-	/* vector functions */
-	T *resize(size_t width, size_t height = 0, size_t depth = 0)
+	/* Host memory allocation. */
+	T *alloc(size_t width, size_t height = 0, size_t depth = 0)
 	{
-		data_size = width * ((height == 0)? 1: height) * ((depth == 0)? 1: depth);
-		if(data.resize(data_size) == NULL) {
-			clear();
-			return NULL;
+		size_t new_size = size(width, height, depth);
+
+		if(new_size != data_size) {
+			device_free();
+			host_free(data_pointer, sizeof(T)*data_size);
+			data_pointer = host_alloc(sizeof(T)*new_size);
 		}
+
+		data_size = new_size;
 		data_width = width;
 		data_height = height;
 		data_depth = depth;
-		if(data_size == 0) {
-			data_pointer = 0;
-			return NULL;
+		assert(device_ptr == 0);
+
+		return get_data();
+	}
+
+	/* Host memory resize. Only use this if the original data needs to be
+	 * preserved, it is faster to call alloc() if it can be discarded. */
+	T *resize(size_t width, size_t height = 0, size_t depth = 0)
+	{
+		size_t new_size = size(width, height, depth);
+
+		if(new_size != data_size) {
+			device_ptr new_ptr = host_alloc(sizeof(T)*new_size);
+
+			if(new_size && data_size) {
+				size_t min_size = ((new_size < data_size)? new_size: data_size);
+				memcpy((T*)new_ptr, (T*)data_pointer, sizeof(T)*min_size);
+			}
+
+			device_free();
+			host_free(data_pointer, sizeof(T)*data_size);
+			data_pointer = new_ptr;
 		}
-		data_pointer = (device_ptr)&data[0];
-		return &data[0];
+
+		data_size = new_size;
+		data_width = width;
+		data_height = height;
+		data_depth = depth;
+		assert(device_ptr == 0);
+
+		return get_data();
 	}
 
+	/* Take over data from an existing array. */
 	void steal_data(array<T>& from)
 	{
-		data.steal_data(from);
-		data_size = data.size();
-		data_pointer = (data_size)? (device_ptr)&data[0]: 0;
-		data_width = data_size;
+		device_free();
+		host_free(data_pointer, sizeof(T)*data_size);
+
+		data_size = from.size();
+		data_width = 0;
 		data_height = 0;
 		data_depth = 0;
+		data_pointer = (device_ptr)from.steal_pointer();
+		assert(device_pointer == 0);
 	}
 
-	void clear()
+	/* Free device and host memory. */
+	void free()
 	{
-		data.clear();
-		data_pointer = 0;
+		device_free();
+		host_free(data_pointer, sizeof(T)*data_size);
+
+		data_size = 0;
 		data_width = 0;
 		data_height = 0;
 		data_depth = 0;
-		data_size = 0;
-		device_pointer = 0;
+		data_pointer = 0;
+		assert(device_pointer == 0);
 	}
 
 	size_t size()
 	{
-		return data.size();
+		return data_size;
 	}
 
 	T* get_data()
 	{
-		return &data[0];
+		return (T*)data_pointer;
 	}
 
 	T& operator[](size_t i)
 	{
-		return data[i];
+		assert(i < data_size);
+		return get_data()[i];
 	}
 
-private:
-	array<T> data;
+	void copy_to_device()
+	{
+		device_copy_to();
+	}
+
+	void copy_from_device(int y, int w, int h)
+	{
+		device_copy_from(y, w, h, sizeof(T));
+	}
+
+	void zero_to_device()
+	{
+		device_zero();
+	}
+
+protected:
+	size_t size(size_t width, size_t height, size_t depth)
+	{
+		return width * ((height == 0)? 1: height) * ((depth == 0)? 1: depth);
+	}
 };
 
-/* A device_sub_ptr is a pointer into another existing memory.
- * Therefore, it is not allocated separately, but just created from the already allocated base memory.
- * It is freed automatically when it goes out of scope, which should happen before the base memory is freed.
- * Note that some devices require the offset and size of the sub_ptr to be properly aligned. */
+/* Pixel Memory
+ *
+ * Device memory to efficiently draw as pixels to the screen in interactive
+ * rendering. Only copying pixels from the device is supported, not copying to. */
+
+template<typename T> class device_pixels : public device_vector<T>
+{
+public:
+	device_pixels(Device *device, const char *name)
+	: device_vector<T>(device, name, MEM_PIXELS)
+	{
+	}
+
+	void alloc_to_device(size_t width, size_t height, size_t depth = 0)
+	{
+		device_vector<T>::alloc(width, height, depth);
+		device_memory::device_alloc();
+	}
+
+	T *copy_from_device(int y, int w, int h)
+	{
+		device_memory::device_copy_from(y, w, h, sizeof(T));
+		return device_vector<T>::get_data();
+	}
+};
+
+/* Device Sub Memory
+ *
+ * Pointer into existing memory. It is not allocated separately, but created
+ * from an already allocated base memory. It is freed automatically when it
+ * goes out of scope, which should happen before base memory is freed.
+ *
+ * Note: some devices require offset and size of the sub_ptr to be properly
+ * aligned to device->mem_address_alingment(). */
+
 class device_sub_ptr
 {
 public:
 	device_sub_ptr(device_memory& mem, int offset, int size);
 	~device_sub_ptr();
-	/* No copying. */
-	device_sub_ptr& operator = (const device_sub_ptr&);
 
 	device_ptr operator*() const
 	{
 		return ptr;
 	}
+
 protected:
+	/* No copying. */
+	device_sub_ptr& operator = (const device_sub_ptr&);
+
 	Device *device;
 	device_ptr ptr;
 };
diff --git a/intern/cycles/device/device_multi.cpp b/intern/cycles/device/device_multi.cpp
index 7f7fbc0d1d3..0a6dd90c86d 100644
--- a/intern/cycles/device/device_multi.cpp
+++ b/intern/cycles/device/device_multi.cpp
@@ -43,10 +43,10 @@ public:
 	};
 
 	list<SubDevice> devices;
-	device_ptr unique_ptr;
+	device_ptr unique_key;
 
 	MultiDevice(DeviceInfo& info, Stats &stats, bool background_)
-	: Device(info, stats, background_), unique_ptr(1)
+	: Device(info, stats, background_), unique_key(1)
 	{
 		Device *device;
 
@@ -108,68 +108,87 @@ public:
 
 	void mem_alloc(device_memory& mem)
 	{
+		device_ptr key = unique_key++;
+
 		foreach(SubDevice& sub, devices) {
+			mem.device = sub.device;
 			mem.device_pointer = 0;
+
 			sub.device->mem_alloc(mem);
-			sub.ptr_map[unique_ptr] = mem.device_pointer;
+			sub.ptr_map[key] = mem.device_pointer;
 		}
 
-		mem.device_pointer = unique_ptr++;
-		stats.mem_alloc(mem.device_size);
+		mem.device = this;
+		mem.device_pointer = key;
 	}
 
 	void mem_copy_to(device_memory& mem)
 	{
-		device_ptr tmp = mem.device_pointer;
+		device_ptr existing_key = mem.device_pointer;
+		device_ptr key = (existing_key)? existing_key: unique_key++;
 
 		foreach(SubDevice& sub, devices) {
-			mem.device_pointer = sub.ptr_map[tmp];
+			mem.device = sub.device;
+			mem.device_pointer = (existing_key)? sub.ptr_map[existing_key]: 0;
+
 			sub.device->mem_copy_to(mem);
+			sub.ptr_map[key] = mem.device_pointer;
 		}
 
-		mem.device_pointer = tmp;
+		mem.device = this;
+		mem.device_pointer = key;
 	}
 
 	void mem_copy_from(device_memory& mem, int y, int w, int h, int elem)
 	{
-		device_ptr tmp = mem.device_pointer;
+		device_ptr key = mem.device_pointer;
 		int i = 0, sub_h = h/devices.size();
 
 		foreach(SubDevice& sub, devices) {
 			int sy = y + i*sub_h;
 			int sh = (i == (int)devices.size() - 1)? h - sub_h*i: sub_h;
 
-			mem.device_pointer = sub.ptr_map[tmp];
+			mem.device = sub.device;
+			mem.device_pointer = sub.ptr_map[key];
+
 			sub.device->mem_copy_from(mem, sy, w, sh, elem);
 			i++;
 		}
 
-		mem.device_pointer = tmp;
+		mem.device = this;
+		mem.device_pointer = key;
 	}
 
 	void mem_zero(device_memory& mem)
 	{
-		device_ptr tmp = mem.device_pointer;
+		device_ptr existing_key = mem.device_pointer;
+		device_ptr key = (existing_key)? existing_key: unique_key++;
 
 		foreach(SubDevice& sub, devices) {
-			mem.device_pointer = sub.ptr_map[tmp];
+			mem.device = sub.device;
+			mem.device_pointer = (existing_key)? sub.ptr_map[existing_key]: 0;
+
 			sub.device->mem_zero(mem);
+			sub.ptr_map[key] = mem.device_pointer;
 		}
 
-		mem.device_pointer = tmp;
+		mem.device = this;
+		mem.device_pointer = key;
 	}
 
 	void mem_free(device_memory& mem)
 	{
-		device_ptr tmp = mem.device_pointer;
-		stats.mem_free(mem.device_size);
+		device_ptr key = mem.device_pointer;
 
 		foreach(SubDevice& sub, devices) {
-			mem.device_pointer = sub.ptr_map[tmp];
+			mem.device = sub.device;
+			mem.device_pointer = sub.ptr_map[key];
+
 			sub.device->mem_free(mem);
-			sub.ptr_map.erase(sub.ptr_map.find(tmp));
+			sub.ptr_map.erase(sub.ptr_map.find(key));
 		}
 
+		mem.device = this;
 		mem.device_pointer = 0;
 	}
 
@@ -179,81 +198,10 @@ public:
 			sub.device->const_copy_to(name, host, size);
 	}
 
-	void tex_alloc(device_memory& mem)
-	{
-		VLOG(1) << "Texture allocate: " << mem.name << ", "
-		        << string_human_readable_number(mem.memory_size()) << " bytes. ("
-		        << string_human_readable_size(mem.memory_size()) << ")";
-
-		foreach(SubDevice& sub, devices) {
-			mem.device_pointer = 0;
-			sub.device->tex_alloc(mem);
-			sub.ptr_map[unique_ptr] = mem.device_pointer;
-		}
-
-		mem.device_pointer = unique_ptr++;
-		stats.mem_alloc(mem.device_size);
-	}
-
-	void tex_free(device_memory& mem)
-	{
-		device_ptr tmp = mem.device_pointer;
-		stats.mem_free(mem.device_size);
-
-		foreach(SubDevice& sub, devices) {
-			mem.device_pointer = sub.ptr_map[tmp];
-			sub.device->tex_free(mem);
-			sub.ptr_map.erase(sub.ptr_map.find(tmp));
-		}
-
-		mem.device_pointer = 0;
-	}
-
-	void pixels_alloc(device_memory& mem)
-	{
-		foreach(SubDevice& sub, devices) {
-			mem.device_pointer = 0;
-			sub.device->pixels_alloc(mem);
-			sub.ptr_map[unique_ptr] = mem.device_pointer;
-		}
-
-		mem.device_pointer = unique_ptr++;
-	}
-
-	void pixels_free(device_memory& mem)
-	{
-		device_ptr tmp = mem.device_pointer;
-
-		foreach(SubDevice& sub, devices) {
-			mem.device_pointer = sub.ptr_map[tmp];
-			sub.device->pixels_free(mem);
-			sub.ptr_map.erase(sub.ptr_map.find(tmp));
-		}
-
-		mem.device_pointer = 0;
-	}
-
-	void pixels_copy_from(device_memory& mem, int y, int w, int h)
-	{
-		device_ptr tmp = mem.device_pointer;
-		int i = 0, sub_h = h/devices.size();
-
-		foreach(SubDevice& sub, devices) {
-			int sy = y + i*sub_h;
-			int sh = (i == (int)devices.size() - 1)? h - sub_h*i: sub_h;
-
-			mem.device_pointer = sub.ptr_map[tmp];
-			sub.device->pixels_copy_from(mem, sy, w, sh);
-			i++;
-		}
-
-		mem.device_pointer = tmp;
-	}
-
 	void draw_pixels(device_memory& rgba, int y, int w, int h, int dx, int dy, int width, int height, bool transparent,
 		const DeviceDrawParams &draw_params)
 	{
-		device_ptr tmp = rgba.device_pointer;
+		device_ptr key = rgba.device_pointer;
 		int i = 0, sub_h = h/devices.size();
 		int sub_height = height/devices.size();
 
@@ -264,12 +212,12 @@ public:
 			int sdy = dy + i*sub_height;
 			/* adjust math for w/width */
 
-			rgba.device_pointer = sub.ptr_map[tmp];
+			rgba.device_pointer = sub.ptr_map[key];
 			sub.device->draw_pixels(rgba, sy, w, sh, dx, sdy, width, sheight, transparent, draw_params);
 			i++;
 		}
 
-		rgba.device_pointer = tmp;
+		rgba.device_pointer = key;
 	}
 
 	void map_tile(Device *sub_device, RenderTile& tile)
@@ -304,15 +252,21 @@ public:
 			 * to the current device now, for the duration of the denoising task.
 			 * Note that this temporarily modifies the RenderBuffers and calls
 			 * the device, so this function is not thread safe. */
-			if(tiles[i].buffers->device != sub_device) {
-				device_vector<float> &mem = tiles[i].buffers->buffer;
-
+			device_vector<float> &mem = tiles[i].buffers->buffer;
+			if(mem.device != sub_device) {
 				tiles[i].buffers->copy_from_device();
+
+				Device *original_device = mem.device;
 				device_ptr original_ptr = mem.device_pointer;
+
+				mem.device = sub_device;
 				mem.device_pointer = 0;
+
 				sub_device->mem_alloc(mem);
 				sub_device->mem_copy_to(mem);
 				tiles[i].buffer = mem.device_pointer;
+
+				mem.device = original_device;
 				mem.device_pointer = original_ptr;
 			}
 		}
@@ -324,25 +278,30 @@ public:
 			if(!tiles[i].buffers) {
 				continue;
 			}
-			if(tiles[i].buffers->device != sub_device) {
-				device_vector<float> &mem = tiles[i].buffers->buffer;
 
+			device_vector<float> &mem = tiles[i].buffers->buffer;
+			if(mem.device != sub_device) {
+				Device *original_device = mem.device;
 				device_ptr original_ptr = mem.device_pointer;
+				size_t original_size = mem.device_size;
+
+				mem.device = sub_device;
 				mem.device_pointer = tiles[i].buffer;
 
 				/* Copy denoised tile to the host. */
 				if(i == 4) {
-					tiles[i].buffers->copy_from_device(sub_device);
+					tiles[i].buffers->copy_from_device();
 				}
 
-				size_t mem_size = mem.device_size;
 				sub_device->mem_free(mem);
+
+				mem.device = original_device;
 				mem.device_pointer = original_ptr;
-				mem.device_size = mem_size;
+				mem.device_size = original_size;
 
 				/* Copy denoised tile to the original device. */
 				if(i == 4) {
-					tiles[i].buffers->device->mem_copy_to(mem);
+					mem.copy_to_device();
 				}
 			}
 		}
diff --git a/intern/cycles/device/device_network.cpp b/intern/cycles/device/device_network.cpp
index bdc88b6acae..fa231c817e6 100644
--- a/intern/cycles/device/device_network.cpp
+++ b/intern/cycles/device/device_network.cpp
@@ -172,36 +172,6 @@ public:
 		snd.write_buffer(host, size);
 	}
 
-	void tex_alloc(device_memory& mem)
-	{
-		VLOG(1) << "Texture allocate: " << mem.name << ", "
-		        << string_human_readable_number(mem.memory_size()) << " bytes. ("
-		        << string_human_readable_size(mem.memory_size()) << ")";
-
-		thread_scoped_lock lock(rpc_lock);
-
-		mem.device_pointer = ++mem_counter;
-
-		RPCSend snd(socket, &error_func, "tex_alloc");
-		snd.add(mem);
-		snd.write();
-		snd.write_buffer((void*)mem.data_pointer, mem.memory_size());
-	}
-
-	void tex_free(device_memory& mem)
-	{
-		if(mem.device_pointer) {
-			thread_scoped_lock lock(rpc_lock);
-
-			RPCSend snd(socket, &error_func, "tex_free");
-
-			snd.add(mem);
-			snd.write();
-
-			mem.device_pointer = 0;
-		}
-	}
-
 	bool load_kernels(const DeviceRequestedFeatures& requested_features)
 	{
 		if(error_func.have_error())
@@ -310,7 +280,7 @@ public:
 		snd.write();
 	}
 
-	int get_split_task_count(DeviceTask& task)
+	int get_split_task_count(DeviceTask&)
 	{
 		return 1;
 	}
@@ -464,21 +434,17 @@ protected:
 			rcv.read(mem, name);
 			lock.unlock();
 
+			/* Allocate host side data buffer. */
+			size_t data_size = mem.memory_size();
 			device_ptr client_pointer = mem.device_pointer;
 
-			/* create a memory buffer for the device buffer */
-			size_t data_size = mem.memory_size();
 			DataVector &data_v = data_vector_insert(client_pointer, data_size);
+			mem.data_pointer = (data_size)? (device_ptr)&(data_v[0]): 0;
 
-			if(data_size)
-				mem.data_pointer = (device_ptr)&(data_v[0]);
-			else
-				mem.data_pointer = 0;
-
-			/* perform the allocation on the actual device */
+			/* Perform the allocation on the actual device. */
 			device->mem_alloc(mem);
 
-			/* store a mapping to/from client_pointer and real device pointer */
+			/* Store a mapping to/from client_pointer and real device pointer. */
 			pointer_mapping_insert(client_pointer, mem.device_pointer);
 		}
 		else if(rcv.name == "mem_copy_to") {
@@ -487,23 +453,33 @@ protected:
 			rcv.read(mem, name);
 			lock.unlock();
 
+			size_t data_size = mem.memory_size();
 			device_ptr client_pointer = mem.device_pointer;
 
-			DataVector &data_v = data_vector_find(client_pointer);
-
-			size_t data_size = mem.memory_size();
+			if(client_pointer) {
+				/* Lookup existing host side data buffer. */
+				DataVector &data_v = data_vector_find(client_pointer);
+				mem.data_pointer = (device_ptr)&data_v[0];
 
-			/* get pointer to memory buffer	for device buffer */
-			mem.data_pointer = (device_ptr)&data_v[0];
+				/* Translate the client pointer to a real device pointer. */
+				mem.device_pointer = device_ptr_from_client_pointer(client_pointer);
+			}
+			else {
+				/* Allocate host side data buffer. */
+				DataVector &data_v = data_vector_insert(client_pointer, data_size);
+				mem.data_pointer = (data_size)? (device_ptr)&(data_v[0]): 0;
+			}
 
-			/* copy data from network into memory buffer */
+			/* Copy data from network into memory buffer. */
 			rcv.read_buffer((uint8_t*)mem.data_pointer, data_size);
 
-			/* translate the client pointer to a real device pointer */
-			mem.device_pointer = device_ptr_from_client_pointer(client_pointer);
-
-			/* copy the data from the memory buffer to the device buffer */
+			/* Copy the data from the memory buffer to the device buffer. */
 			device->mem_copy_to(mem);
+
+			if(!client_pointer) {
+				/* Store a mapping to/from client_pointer and real device pointer. */
+				pointer_mapping_insert(client_pointer, mem.device_pointer);
+			}
 		}
 		else if(rcv.name == "mem_copy_from") {
 			string name;
@@ -538,14 +514,30 @@ protected:
 			rcv.read(mem, name);
 			lock.unlock();
 
+			size_t data_size = mem.memory_size();
 			device_ptr client_pointer = mem.device_pointer;
-			mem.device_pointer = device_ptr_from_client_pointer(client_pointer);
 
-			DataVector &data_v = data_vector_find(client_pointer);
+			if(client_pointer) {
+				/* Lookup existing host side data buffer. */
+				DataVector &data_v = data_vector_find(client_pointer);
+				mem.data_pointer = (device_ptr)&data_v[0];
 
-			mem.data_pointer = (device_ptr)&(data_v[0]);
+				/* Translate the client pointer to a real device pointer. */
+				mem.device_pointer = device_ptr_from_client_pointer(client_pointer);
+			}
+			else {
+				/* Allocate host side data buffer. */
+				DataVector &data_v = data_vector_insert(client_pointer, data_size);
+				mem.data_pointer = (data_size)? (device_ptr)&(data_v[0]): 0;
+			}
 
+			/* Zero memory. */
 			device->mem_zero(mem);
+
+			if(!client_pointer) {
+				/* Store a mapping to/from client_pointer and real device pointer. */
+				pointer_mapping_insert(client_pointer, mem.device_pointer);
+			}
 		}
 		else if(rcv.name == "mem_free") {
 			string name;
@@ -573,45 +565,6 @@ protected:
 
 			device->const_copy_to(name_string.c_str(), &host_vector[0], size);
 		}
-		else if(rcv.name == "tex_alloc") {
-			string name;
-			network_device_memory mem(device);
-			device_ptr client_pointer;
-
-			rcv.read(mem, name);
-			lock.unlock();
-
-			client_pointer = mem.device_pointer;
-
-			size_t data_size = mem.memory_size();
-
-			DataVector &data_v = data_vector_insert(client_pointer, data_size);
-
-			if(data_size)
-				mem.data_pointer = (device_ptr)&(data_v[0]);
-			else
-				mem.data_pointer = 0;
-
-			rcv.read_buffer((uint8_t*)mem.data_pointer, data_size);
-
-			device->tex_alloc(mem);
-
-			pointer_mapping_insert(client_pointer, mem.device_pointer);
-		}
-		else if(rcv.name == "tex_free") {
-			string name;
-			network_device_memory mem(device);
-			device_ptr client_pointer;
-
-			rcv.read(mem, name);
-			lock.unlock();
-
-			client_pointer = mem.device_pointer;
-
-			mem.device_pointer = device_ptr_from_client_pointer_erase(client_pointer);
-
-			device->tex_free(mem);
-		}
 		else if(rcv.name == "load_kernels") {
 			DeviceRequestedFeatures requested_features;
 			rcv.read(requested_features.experimental);
@@ -696,7 +649,7 @@ protected:
 		}
 	}
 
-	bool task_acquire_tile(Device *device, RenderTile& tile)
+	bool task_acquire_tile(Device *, RenderTile& tile)
 	{
 		thread_scoped_lock acquire_lock(acquire_mutex);
 
diff --git a/intern/cycles/device/device_network.h b/intern/cycles/device/device_network.h
index 8a53290f421..a38d962c0af 100644
--- a/intern/cycles/device/device_network.h
+++ b/intern/cycles/device/device_network.h
@@ -279,6 +279,11 @@ public:
 
 		mem.name = name.c_str();
 		mem.data_pointer = 0;
+
+		/* Can't transfer OpenGL texture over network. */
+		if(mem.type == MEM_PIXELS) {
+			mem.type = MEM_WRITE_ONLY;
+		}
 	}
 
 	template<typename T> void read(T& data)
diff --git a/intern/cycles/device/device_split_kernel.cpp b/intern/cycles/device/device_split_kernel.cpp
index 6c8befa89be..f2839a8b1b9 100644
--- a/intern/cycles/device/device_split_kernel.cpp
+++ b/intern/cycles/device/device_split_kernel.cpp
@@ -61,11 +61,11 @@ DeviceSplitKernel::DeviceSplitKernel(Device *device)
 
 DeviceSplitKernel::~DeviceSplitKernel()
 {
-	device->mem_free(split_data);
-	device->mem_free(ray_state);
-	device->mem_free(use_queues_flag);
-	device->mem_free(queue_index);
-	device->mem_free(work_pool_wgs);
+	split_data.free();
+	ray_state.free();
+	use_queues_flag.free();
+	queue_index.free();
+	work_pool_wgs.free();
 
 	delete kernel_path_init;
 	delete kernel_scene_intersect;
@@ -175,20 +175,11 @@ bool DeviceSplitKernel::path_trace(DeviceTask *task,
 		unsigned int max_work_groups = num_global_elements / work_pool_size + 1;
 
 		/* Allocate work_pool_wgs memory. */
-		work_pool_wgs.resize(max_work_groups);
-		device->mem_alloc(work_pool_wgs);
-
-		queue_index.resize(NUM_QUEUES);
-		device->mem_alloc(queue_index);
-
-		use_queues_flag.resize(1);
-		device->mem_alloc(use_queues_flag);
-
-		ray_state.resize(num_global_elements);
-		device->mem_alloc(ray_state);
-
-		split_data.resize(state_buffer_size(kgbuffer, kernel_data, num_global_elements));
-		device->mem_alloc(split_data);
+		work_pool_wgs.alloc_to_device(max_work_groups);
+		queue_index.alloc_to_device(NUM_QUEUES);
+		use_queues_flag.alloc_to_device(1);
+		split_data.alloc_to_device(state_buffer_size(kgbuffer, kernel_data, num_global_elements));
+		ray_state.alloc(num_global_elements);
 	}
 
 #define ENQUEUE_SPLIT_KERNEL(name, global_size, local_size) \
@@ -225,9 +216,9 @@ bool DeviceSplitKernel::path_trace(DeviceTask *task,
 		/* reset state memory here as global size for data_init
 		 * kernel might not be large enough to do in kernel
 		 */
-		device->mem_zero(work_pool_wgs);
-		device->mem_zero(split_data);
-		device->mem_zero(ray_state);
+		work_pool_wgs.zero_to_device();
+		split_data.zero_to_device();
+		ray_state.zero_to_device();
 
 		if(!enqueue_split_kernel_data_init(KernelDimensions(global_size, local_size),
 		                                   subtile,
@@ -284,7 +275,7 @@ bool DeviceSplitKernel::path_trace(DeviceTask *task,
 			}
 
 			/* Decide if we should exit path-iteration in host. */
-			device->mem_copy_from(ray_state, 0, global_size[0] * global_size[1] * sizeof(char), 1, 1);
+			ray_state.copy_from_device(0, global_size[0] * global_size[1], 1);
 
 			activeRaysAvailable = false;
 
diff --git a/intern/cycles/device/opencl/memory_manager.cpp b/intern/cycles/device/opencl/memory_manager.cpp
index e48367b8987..a791b374774 100644
--- a/intern/cycles/device/opencl/memory_manager.cpp
+++ b/intern/cycles/device/opencl/memory_manager.cpp
@@ -76,8 +76,7 @@ void MemoryManager::DeviceBuffer::update_device_memory(OpenCLDeviceBase *device)
 		device_only_memory<uchar> *new_buffer =
 			new device_only_memory<uchar>(device, "memory manager buffer");
 
-		new_buffer->resize(total_size);
-		device->mem_alloc(*new_buffer);
+		new_buffer->alloc_to_device(total_size);
 
 		size_t offset = 0;
 
@@ -111,7 +110,6 @@ void MemoryManager::DeviceBuffer::update_device_memory(OpenCLDeviceBase *device)
 			offset += allocation->size;
 		}
 
-		device->mem_free(*buffer);
 		delete buffer;
 
 		buffer = new_buffer;
@@ -144,9 +142,9 @@ void MemoryManager::DeviceBuffer::update_device_memory(OpenCLDeviceBase *device)
 	clFinish(device->cqCommandQueue);
 }
 
-void MemoryManager::DeviceBuffer::free(OpenCLDeviceBase *device)
+void MemoryManager::DeviceBuffer::free(OpenCLDeviceBase *)
 {
-	device->mem_free(*buffer);
+	buffer->free();
 }
 
 MemoryManager::DeviceBuffer* MemoryManager::smallest_device_buffer()
diff --git a/intern/cycles/device/opencl/opencl_base.cpp b/intern/cycles/device/opencl/opencl_base.cpp
index 90f461b4c98..5e9debc3b17 100644
--- a/intern/cycles/device/opencl/opencl_base.cpp
+++ b/intern/cycles/device/opencl/opencl_base.cpp
@@ -74,7 +74,7 @@ void OpenCLDeviceBase::opencl_assert_err(cl_int err, const char* where)
 OpenCLDeviceBase::OpenCLDeviceBase(DeviceInfo& info, Stats &stats, bool background_)
 : Device(info, stats, background_),
   memory_manager(this),
-  texture_info_buffer(this, "__texture_info", MEM_READ_ONLY)
+  texture_info(this, "__texture_info", MEM_TEXTURE)
 {
 	cpPlatform = NULL;
 	cdDevice = NULL;
@@ -157,7 +157,6 @@ OpenCLDeviceBase::~OpenCLDeviceBase()
 
 	ConstMemMap::iterator mt;
 	for(mt = const_mem_map.begin(); mt != const_mem_map.end(); mt++) {
-		mem_free(*(mt->second));
 		delete mt->second;
 	}
 
@@ -318,9 +317,9 @@ void OpenCLDeviceBase::mem_alloc(device_memory& mem)
 	cl_mem_flags mem_flag;
 	void *mem_ptr = NULL;
 
-	if(mem.type == MEM_READ_ONLY)
+	if(mem.type == MEM_READ_ONLY || mem.type == MEM_TEXTURE)
 		mem_flag = CL_MEM_READ_ONLY;
-	else if(mem.type == MEM_WRITE_ONLY)
+	else if(mem.type == MEM_WRITE_ONLY || mem.type == MEM_PIXELS)
 		mem_flag = CL_MEM_WRITE_ONLY;
 	else
 		mem_flag = CL_MEM_READ_WRITE;
@@ -348,17 +347,27 @@ void OpenCLDeviceBase::mem_alloc(device_memory& mem)
 
 void OpenCLDeviceBase::mem_copy_to(device_memory& mem)
 {
-	/* this is blocking */
-	size_t size = mem.memory_size();
-	if(size != 0) {
-		opencl_assert(clEnqueueWriteBuffer(cqCommandQueue,
-		                                   CL_MEM_PTR(mem.device_pointer),
-		                                   CL_TRUE,
-		                                   0,
-		                                   size,
-		                                   (void*)mem.data_pointer,
-		                                   0,
-		                                   NULL, NULL));
+	if(mem.type == MEM_TEXTURE) {
+		tex_free(mem);
+		tex_alloc(mem);
+	}
+	else {
+		if(!mem.device_pointer) {
+			mem_alloc(mem);
+		}
+
+		/* this is blocking */
+		size_t size = mem.memory_size();
+		if(size != 0) {
+			opencl_assert(clEnqueueWriteBuffer(cqCommandQueue,
+			                                   CL_MEM_PTR(mem.device_pointer),
+			                                   CL_TRUE,
+			                                   0,
+			                                   size,
+			                                   (void*)mem.data_pointer,
+			                                   0,
+			                                   NULL, NULL));
+		}
 	}
 }
 
@@ -410,6 +419,10 @@ void OpenCLDeviceBase::mem_zero_kernel(device_ptr mem, size_t size)
 
 void OpenCLDeviceBase::mem_zero(device_memory& mem)
 {
+	if(!mem.device_pointer) {
+		mem_alloc(mem);
+	}
+
 	if(mem.device_pointer) {
 		if(base_program.is_loaded()) {
 			mem_zero_kernel(mem.device_pointer, mem.memory_size());
@@ -445,14 +458,19 @@ void OpenCLDeviceBase::mem_zero(device_memory& mem)
 
 void OpenCLDeviceBase::mem_free(device_memory& mem)
 {
-	if(mem.device_pointer) {
-		if(mem.device_pointer != null_mem) {
-			opencl_assert(clReleaseMemObject(CL_MEM_PTR(mem.device_pointer)));
-		}
-		mem.device_pointer = 0;
+	if(mem.type == MEM_TEXTURE) {
+		tex_free(mem);
+	}
+	else {
+		if(mem.device_pointer) {
+			if(mem.device_pointer != null_mem) {
+				opencl_assert(clReleaseMemObject(CL_MEM_PTR(mem.device_pointer)));
+			}
+			mem.device_pointer = 0;
 
-		stats.mem_free(mem.device_size);
-		mem.device_size = 0;
+			stats.mem_free(mem.device_size);
+			mem.device_size = 0;
+		}
 	}
 }
 
@@ -464,9 +482,9 @@ int OpenCLDeviceBase::mem_address_alignment()
 device_ptr OpenCLDeviceBase::mem_alloc_sub_ptr(device_memory& mem, int offset, int size)
 {
 	cl_mem_flags mem_flag;
-	if(mem.type == MEM_READ_ONLY)
+	if(mem.type == MEM_READ_ONLY || mem.type == MEM_TEXTURE)
 		mem_flag = CL_MEM_READ_ONLY;
-	else if(mem.type == MEM_WRITE_ONLY)
+	else if(mem.type == MEM_WRITE_ONLY || mem.type == MEM_PIXELS)
 		mem_flag = CL_MEM_WRITE_ONLY;
 	else
 		mem_flag = CL_MEM_READ_WRITE;
@@ -498,9 +516,7 @@ void OpenCLDeviceBase::const_copy_to(const char *name, void *host, size_t size)
 
 	if(i == const_mem_map.end()) {
 		data = new device_vector<uchar>(this, name, MEM_READ_ONLY);
-		data->resize(size);
-
-		mem_alloc(*data);
+		data->alloc(size);
 		const_mem_map.insert(ConstMemMap::value_type(name, data));
 	}
 	else {
@@ -508,7 +524,7 @@ void OpenCLDeviceBase::const_copy_to(const char *name, void *host, size_t size)
 	}
 
 	memcpy(data->get_data(), host, size);
-	mem_copy_to(*data);
+	data->copy_to_device();
 }
 
 void OpenCLDeviceBase::tex_alloc(device_memory& mem)
@@ -1037,8 +1053,7 @@ bool OpenCLDeviceBase::denoising_detect_outliers(device_ptr image_ptr,
 bool OpenCLDeviceBase::denoising_set_tiles(device_ptr *buffers,
                                            DenoisingTask *task)
 {
-	mem_alloc(task->tiles_mem);
-	mem_copy_to(task->tiles_mem);
+	task->tiles_mem.copy_to_device();
 
 	cl_mem tiles_mem = CL_MEM_PTR(task->tiles_mem.device_pointer);
 
diff --git a/intern/cycles/device/opencl/opencl_split.cpp b/intern/cycles/device/opencl/opencl_split.cpp
index c966ebe0c5e..96139afa450 100644
--- a/intern/cycles/device/opencl/opencl_split.cpp
+++ b/intern/cycles/device/opencl/opencl_split.cpp
@@ -128,8 +128,7 @@ public:
 
 			/* Allocate buffer for kernel globals */
 			device_only_memory<KernelGlobals> kgbuffer(this, "kernel_globals");
-			kgbuffer.resize(1);
-			mem_alloc(kgbuffer);
+			kgbuffer.alloc_to_device(1);
 
 			/* Keep rendering tiles until done. */
 			while(task->acquire_tile(this, tile)) {
@@ -160,7 +159,7 @@ public:
 				task->release_tile(tile);
 			}
 
-			mem_free(kgbuffer);
+			kgbuffer.free();
 		}
 	}
 
@@ -289,8 +288,8 @@ public:
 	virtual uint64_t state_buffer_size(device_memory& kg, device_memory& data, size_t num_threads)
 	{
 		device_vector<uint64_t> size_buffer(device, "size_buffer", MEM_READ_WRITE);
-		size_buffer.resize(1);
-		device->mem_alloc(size_buffer);
+		size_buffer.alloc(1);
+		size_buffer.zero_to_device();
 
 		uint threads = num_threads;
 		device->kernel_set_args(device->program_state_buffer_size(), 0, kg, data, threads, size_buffer);
@@ -308,9 +307,9 @@ public:
 
 		device->opencl_assert_err(device->ciErr, "clEnqueueNDRangeKernel");
 
-		device->mem_copy_from(size_buffer, 0, 1, 1, sizeof(uint64_t));
+		size_buffer.copy_from_device(0, 1, 1);
 		size_t size = size_buffer[0];
-		device->mem_free(size_buffer);
+		size_buffer.free();
 
 		if(device->ciErr != CL_SUCCESS) {
 			string message = string_printf("OpenCL error: %s in clEnqueueNDRangeKernel()",
diff --git a/intern/cycles/render/bake.cpp b/intern/cycles/render/bake.cpp
index 66615bf336c..99f68b6aa00 100644
--- a/intern/cycles/render/bake.cpp
+++ b/intern/cycles/render/bake.cpp
@@ -151,7 +151,7 @@ bool BakeManager::bake(Device *device, DeviceScene *dscene, Scene *scene, Progre
 
 		/* setup input for device task */
 		device_vector<uint4> d_input(device, "bake_input", MEM_READ_ONLY);
-		uint4 *d_input_data = d_input.resize(shader_size * 2);
+		uint4 *d_input_data = d_input.alloc(shader_size * 2);
 		size_t d_input_size = 0;
 
 		for(size_t i = shader_offset; i < (shader_offset + shader_size); i++) {
@@ -166,16 +166,13 @@ bool BakeManager::bake(Device *device, DeviceScene *dscene, Scene *scene, Progre
 
 		/* run device task */
 		device_vector<float4> d_output(device, "bake_output", MEM_READ_WRITE);
-		d_output.resize(shader_size);
+		d_output.alloc(shader_size);
+		d_output.zero_to_device();
+		d_input.copy_to_device();
 
 		/* needs to be up to data for attribute access */
 		device->const_copy_to("__data", &dscene->data, sizeof(dscene->data));
 
-		device->mem_alloc(d_input);
-		device->mem_copy_to(d_input);
-		device->mem_alloc(d_output);
-		device->mem_zero(d_output);
-
 		DeviceTask task(DeviceTask::SHADER);
 		task.shader_input = d_input.device_pointer;
 		task.shader_output = d_output.device_pointer;
@@ -192,15 +189,14 @@ bool BakeManager::bake(Device *device, DeviceScene *dscene, Scene *scene, Progre
 		device->task_wait();
 
 		if(progress.get_cancel()) {
-			device->mem_free(d_input);
-			device->mem_free(d_output);
+			d_input.free();
+			d_output.free();
 			m_is_baking = false;
 			return false;
 		}
 
-		device->mem_copy_from(d_output, 0, 1, d_output.size(), sizeof(float4));
-		device->mem_free(d_input);
-		device->mem_free(d_output);
+		d_output.copy_from_device(0, 1, d_output.size());
+		d_input.free();
 
 		/* read result */
 		int k = 0;
@@ -218,6 +214,8 @@ bool BakeManager::bake(Device *device, DeviceScene *dscene, Scene *scene, Progre
 				}
 			}
 		}
+
+		d_output.free();
 	}
 
 	m_is_baking = false;
diff --git a/intern/cycles/render/buffers.cpp b/intern/cycles/render/buffers.cpp
index 2342dd52d86..01f853dda71 100644
--- a/intern/cycles/render/buffers.cpp
+++ b/intern/cycles/render/buffers.cpp
@@ -115,54 +115,35 @@ RenderTile::RenderTile()
 /* Render Buffers */
 
 RenderBuffers::RenderBuffers(Device *device)
-: buffer(device, "RenderBuffers", MEM_READ_WRITE),
-  device(device)
+: buffer(device, "RenderBuffers", MEM_READ_WRITE)
 {
 }
 
 RenderBuffers::~RenderBuffers()
 {
-	device_free();
+	buffer.free();
 }
 
-void RenderBuffers::device_free()
-{
-	if(buffer.device_pointer) {
-		device->mem_free(buffer);
-		buffer.clear();
-	}
-}
-
-void RenderBuffers::reset(Device *device, BufferParams& params_)
+void RenderBuffers::reset(BufferParams& params_)
 {
 	params = params_;
 
-	/* free existing buffers */
-	device_free();
-
-	/* allocate buffer */
-	buffer.resize(params.width*params.height*params.get_passes_size());
-	device->mem_alloc(buffer);
-	device->mem_zero(buffer);
+	/* re-allocate buffer */
+	buffer.alloc(params.width*params.height*params.get_passes_size());
+	buffer.zero_to_device();
 }
 
-void RenderBuffers::zero(Device *device)
+void RenderBuffers::zero()
 {
-	if(buffer.device_pointer) {
-		device->mem_zero(buffer);
-	}
+	buffer.zero_to_device();
 }
 
-bool RenderBuffers::copy_from_device(Device *from_device)
+bool RenderBuffers::copy_from_device()
 {
 	if(!buffer.device_pointer)
 		return false;
 
-	if(!from_device) {
-		from_device = device;
-	}
-
-	from_device->mem_copy_from(buffer, 0, params.width, params.height, params.get_passes_size()*sizeof(float));
+	buffer.copy_from_device(0, params.width * params.get_passes_size(), params.height);
 
 	return true;
 }
@@ -402,47 +383,30 @@ DisplayBuffer::DisplayBuffer(Device *device, bool linear)
   draw_height(0),
   transparent(true), /* todo: determine from background */
   half_float(linear),
-  rgba_byte(device, "display buffer byte", MEM_WRITE_ONLY),
-  rgba_half(device, "display buffer half", MEM_WRITE_ONLY),
-  device(device)
+  rgba_byte(device, "display buffer byte"),
+  rgba_half(device, "display buffer half")
 {
 }
 
 DisplayBuffer::~DisplayBuffer()
 {
-	device_free();
-}
-
-void DisplayBuffer::device_free()
-{
-	if(rgba_byte.device_pointer) {
-		device->pixels_free(rgba_byte);
-		rgba_byte.clear();
-	}
-	if(rgba_half.device_pointer) {
-		device->pixels_free(rgba_half);
-		rgba_half.clear();
-	}
+	rgba_byte.free();
+	rgba_half.free();
 }
 
-void DisplayBuffer::reset(Device *device, BufferParams& params_)
+void DisplayBuffer::reset(BufferParams& params_)
 {
 	draw_width = 0;
 	draw_height = 0;
 
 	params = params_;
 
-	/* free existing buffers */
-	device_free();
-
 	/* allocate display pixels */
 	if(half_float) {
-		rgba_half.resize(params.width, params.height);
-		device->pixels_alloc(rgba_half);
+		rgba_half.alloc_to_device(params.width, params.height);
 	}
 	else {
-		rgba_byte.resize(params.width, params.height);
-		device->pixels_alloc(rgba_byte);
+		rgba_byte.alloc_to_device(params.width, params.height);
 	}
 }
 
@@ -457,7 +421,8 @@ void DisplayBuffer::draw_set(int width, int height)
 void DisplayBuffer::draw(Device *device, const DeviceDrawParams& draw_params)
 {
 	if(draw_width != 0 && draw_height != 0) {
-		device_memory& rgba = rgba_data();
+		device_memory& rgba = (half_float)? (device_memory&)rgba_half:
+		                                    (device_memory&)rgba_byte;
 
 		device->draw_pixels(rgba, 0, draw_width, draw_height, params.full_x, params.full_y, params.width, params.height, transparent, draw_params);
 	}
@@ -468,7 +433,7 @@ bool DisplayBuffer::draw_ready()
 	return (draw_width != 0 && draw_height != 0);
 }
 
-void DisplayBuffer::write(Device *device, const string& filename)
+void DisplayBuffer::write(const string& filename)
 {
 	int w = draw_width;
 	int h = draw_height;
@@ -480,21 +445,19 @@ void DisplayBuffer::write(Device *device, const string& filename)
 		return;
 
 	/* read buffer from device */
-	device_memory& rgba = rgba_data();
-	device->pixels_copy_from(rgba, 0, w, h);
+	uchar4 *pixels = rgba_byte.copy_from_device(0, w, h);
 
 	/* write image */
 	ImageOutput *out = ImageOutput::create(filename);
 	ImageSpec spec(w, h, 4, TypeDesc::UINT8);
-	int scanlinesize = w*4*sizeof(uchar);
 
 	out->open(filename, spec);
 
 	/* conversion for different top/bottom convention */
 	out->write_image(TypeDesc::UINT8,
-		(uchar*)rgba.data_pointer + (h-1)*scanlinesize,
+		(uchar*)(pixels + (h-1)*w),
 		AutoStride,
-		-scanlinesize,
+		-w*sizeof(uchar4),
 		AutoStride);
 
 	out->close();
@@ -502,13 +465,5 @@ void DisplayBuffer::write(Device *device, const string& filename)
 	delete out;
 }
 
-device_memory& DisplayBuffer::rgba_data()
-{
-	if(half_float)
-		return rgba_half;
-	else
-		return rgba_byte;
-}
-
 CCL_NAMESPACE_END
 
diff --git a/intern/cycles/render/buffers.h b/intern/cycles/render/buffers.h
index 2780fc8a68d..8563d6674ec 100644
--- a/intern/cycles/render/buffers.h
+++ b/intern/cycles/render/buffers.h
@@ -75,20 +75,15 @@ public:
 	/* float buffer */
 	device_vector<float> buffer;
 
-	Device *device;
-
 	explicit RenderBuffers(Device *device);
 	~RenderBuffers();
 
-	void reset(Device *device, BufferParams& params);
-	void zero(Device *device);
+	void reset(BufferParams& params);
+	void zero();
 
-	bool copy_from_device(Device *from_device = NULL);
+	bool copy_from_device();
 	bool get_pass_rect(PassType type, float exposure, int sample, int components, float *pixels);
 	bool get_denoising_pass_rect(int offset, float exposure, int sample, int components, float *pixels);
-
-protected:
-	void device_free();
 };
 
 /* Display Buffer
@@ -109,25 +104,18 @@ public:
 	/* use half float? */
 	bool half_float;
 	/* byte buffer for converted result */
-	device_vector<uchar4> rgba_byte;
-	device_vector<half4> rgba_half;
+	device_pixels<uchar4> rgba_byte;
+	device_pixels<half4> rgba_half;
 
 	DisplayBuffer(Device *device, bool linear = false);
 	~DisplayBuffer();
 
-	void reset(Device *device, BufferParams& params);
-	void write(Device *device, const string& filename);
+	void reset(BufferParams& params);
+	void write(const string& filename);
 
 	void draw_set(int width, int height);
 	void draw(Device *device, const DeviceDrawParams& draw_params);
 	bool draw_ready();
-
-	device_memory& rgba_data();
-
-protected:
-	void device_free();
-
-	Device *device;
 };
 
 /* Render Tile
diff --git a/intern/cycles/render/image.cpp b/intern/cycles/render/image.cpp
index e7f5ff002b7..625901ff258 100644
--- a/intern/cycles/render/image.cpp
+++ b/intern/cycles/render/image.cpp
@@ -532,7 +532,8 @@ bool ImageManager::file_load_image(Image *img,
 		pixels = &pixels_storage[0];
 	}
 	else {
-		pixels = (StorageType*)tex_img.resize(width, height, depth);
+		thread_scoped_lock device_lock(device_mutex);
+		pixels = (StorageType*)tex_img.alloc(width, height, depth);
 	}
 	if(pixels == NULL) {
 		/* Could be that we've run out of memory. */
@@ -686,9 +687,16 @@ bool ImageManager::file_load_image(Image *img,
 		                         scale_factor,
 		                         &scaled_pixels,
 		                         &scaled_width, &scaled_height, &scaled_depth);
-		StorageType *texture_pixels = (StorageType*)tex_img.resize(scaled_width,
-		                                                           scaled_height,
-		                                                           scaled_depth);
+
+		StorageType *texture_pixels;
+
+		{
+			thread_scoped_lock device_lock(device_mutex);
+			texture_pixels = (StorageType*)tex_img.alloc(scaled_width,
+			                                             scaled_height,
+			                                             scaled_depth);
+		}
+
 		memcpy(texture_pixels,
 		       &scaled_pixels[0],
 		       scaled_pixels.size() * sizeof(StorageType));
@@ -722,14 +730,14 @@ void ImageManager::device_load_image(Device *device,
 	/* 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;
 	}
 
 	/* Create new texture. */
 	if(type == IMAGE_DATA_TYPE_FLOAT4) {
-		device_vector<float4> *tex_img = new device_vector<float4>(device, name.c_str());
+		device_vector<float4> *tex_img
+			= new device_vector<float4>(device, name.c_str(), MEM_TEXTURE);
 
 		if(!file_load_image<TypeDesc::FLOAT, float>(img,
 		                                            type,
@@ -737,7 +745,7 @@ void ImageManager::device_load_image(Device *device,
 		                                            *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->alloc(1, 1);
 
 			pixels[0] = TEX_IMAGE_MISSING_R;
 			pixels[1] = TEX_IMAGE_MISSING_G;
@@ -746,9 +754,15 @@ void ImageManager::device_load_image(Device *device,
 		}
 
 		img->mem = tex_img;
+		img->mem->interpolation = img->interpolation;
+		img->mem->extension = img->extension;
+
+		thread_scoped_lock device_lock(device_mutex);
+		tex_img->copy_to_device();
 	}
 	else if(type == IMAGE_DATA_TYPE_FLOAT) {
-		device_vector<float> *tex_img = new device_vector<float>(device, name.c_str());
+		device_vector<float> *tex_img
+			= new device_vector<float>(device, name.c_str(), MEM_TEXTURE);
 
 		if(!file_load_image<TypeDesc::FLOAT, float>(img,
 		                                            type,
@@ -756,15 +770,21 @@ void ImageManager::device_load_image(Device *device,
 		                                            *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->alloc(1, 1);
 
 			pixels[0] = TEX_IMAGE_MISSING_R;
 		}
 
 		img->mem = tex_img;
+		img->mem->interpolation = img->interpolation;
+		img->mem->extension = img->extension;
+
+		thread_scoped_lock device_lock(device_mutex);
+		tex_img->copy_to_device();
 	}
 	else if(type == IMAGE_DATA_TYPE_BYTE4) {
-		device_vector<uchar4> *tex_img = new device_vector<uchar4>(device, name.c_str());
+		device_vector<uchar4> *tex_img
+			= new device_vector<uchar4>(device, name.c_str(), MEM_TEXTURE);
 
 		if(!file_load_image<TypeDesc::UINT8, uchar>(img,
 		                                            type,
@@ -772,7 +792,7 @@ void ImageManager::device_load_image(Device *device,
 		                                            *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->alloc(1, 1);
 
 			pixels[0] = (TEX_IMAGE_MISSING_R * 255);
 			pixels[1] = (TEX_IMAGE_MISSING_G * 255);
@@ -781,31 +801,43 @@ void ImageManager::device_load_image(Device *device,
 		}
 
 		img->mem = tex_img;
+		img->mem->interpolation = img->interpolation;
+		img->mem->extension = img->extension;
+
+		thread_scoped_lock device_lock(device_mutex);
+		tex_img->copy_to_device();
 	}
 	else if(type == IMAGE_DATA_TYPE_BYTE) {
-		device_vector<uchar> *tex_img = new device_vector<uchar>(device, name.c_str());
+		device_vector<uchar> *tex_img
+			= new device_vector<uchar>(device, name.c_str(), MEM_TEXTURE);
 
 		if(!file_load_image<TypeDesc::UINT8, uchar>(img,
 		                                            type,
 		                                            texture_limit,
 		                                            *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->alloc(1, 1);
 
 			pixels[0] = (TEX_IMAGE_MISSING_R * 255);
 		}
 
 		img->mem = tex_img;
+		img->mem->interpolation = img->interpolation;
+		img->mem->extension = img->extension;
+
+		thread_scoped_lock device_lock(device_mutex);
+		tex_img->copy_to_device();
 	}
 	else if(type == IMAGE_DATA_TYPE_HALF4) {
-		device_vector<half4> *tex_img = new device_vector<half4>(device, name.c_str());
+		device_vector<half4> *tex_img
+			= new device_vector<half4>(device, name.c_str(), MEM_TEXTURE);
 
 		if(!file_load_image<TypeDesc::HALF, half>(img,
 		                                          type,
 		                                          texture_limit,
 		                                          *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->alloc(1, 1);
 
 			pixels[0] = TEX_IMAGE_MISSING_R;
 			pixels[1] = TEX_IMAGE_MISSING_G;
@@ -814,37 +846,38 @@ void ImageManager::device_load_image(Device *device,
 		}
 
 		img->mem = tex_img;
+		img->mem->interpolation = img->interpolation;
+		img->mem->extension = img->extension;
+
+		thread_scoped_lock device_lock(device_mutex);
+		tex_img->copy_to_device();
 	}
 	else if(type == IMAGE_DATA_TYPE_HALF) {
-		device_vector<half> *tex_img = new device_vector<half>(device, name.c_str());
+		device_vector<half> *tex_img
+			= new device_vector<half>(device, name.c_str(), MEM_TEXTURE);
 
 		if(!file_load_image<TypeDesc::HALF, half>(img,
 		                                          type,
 		                                          texture_limit,
 		                                          *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->alloc(1, 1);
 
 			pixels[0] = TEX_IMAGE_MISSING_R;
 		}
 
 		img->mem = tex_img;
-	}
-
-	/* Copy to device. */
-	if(img->mem) {
 		img->mem->interpolation = img->interpolation;
 		img->mem->extension = img->extension;
 
 		thread_scoped_lock device_lock(device_mutex);
-		device->tex_alloc(*img->mem);
+		tex_img->copy_to_device();
 	}
 
-
 	img->need_load = false;
 }
 
-void ImageManager::device_free_image(Device *device, ImageDataType type, int slot)
+void ImageManager::device_free_image(Device *, ImageDataType type, int slot)
 {
 	Image *img = images[type][slot];
 
@@ -858,7 +891,6 @@ void ImageManager::device_free_image(Device *device, ImageDataType type, int slo
 
 		if(img->mem) {
 			thread_scoped_lock device_lock(device_mutex);
-			device->tex_free(*img->mem);
 			delete img->mem;
 		}
 
diff --git a/intern/cycles/render/integrator.cpp b/intern/cycles/render/integrator.cpp
index b128f18db08..33c3dac9e81 100644
--- a/intern/cycles/render/integrator.cpp
+++ b/intern/cycles/render/integrator.cpp
@@ -191,11 +191,11 @@ void Integrator::device_update(Device *device, DeviceScene *dscene, Scene *scene
 	int dimensions = PRNG_BASE_NUM + max_samples*PRNG_BOUNCE_NUM;
 	dimensions = min(dimensions, SOBOL_MAX_DIMENSIONS);
 
-	uint *directions = dscene->sobol_directions.resize(SOBOL_BITS*dimensions);
+	uint *directions = dscene->sobol_directions.alloc(SOBOL_BITS*dimensions);
 
 	sobol_generate_direction_vectors((uint(*)[SOBOL_BITS])directions, dimensions);
 
-	device->tex_alloc(dscene->sobol_directions);
+	dscene->sobol_directions.copy_to_device();
 
 	/* Clamping. */
 	bool use_sample_clamp = (sample_clamp_direct != 0.0f ||
@@ -208,10 +208,9 @@ void Integrator::device_update(Device *device, DeviceScene *dscene, Scene *scene
 	need_update = false;
 }
 
-void Integrator::device_free(Device *device, DeviceScene *dscene)
+void Integrator::device_free(Device *, DeviceScene *dscene)
 {
-	device->tex_free(dscene->sobol_directions);
-	dscene->sobol_directions.clear();
+	dscene->sobol_directions.free();
 }
 
 bool Integrator::modified(const Integrator& integrator)
diff --git a/intern/cycles/render/light.cpp b/intern/cycles/render/light.cpp
index 9664e1310d5..b3804f34963 100644
--- a/intern/cycles/render/light.cpp
+++ b/intern/cycles/render/light.cpp
@@ -39,7 +39,7 @@ static void shade_background_pixels(Device *device, DeviceScene *dscene, int res
 	device_vector<uint4> d_input(device, "background_input", MEM_READ_ONLY);
 	device_vector<float4> d_output(device, "background_output", MEM_WRITE_ONLY);
 
-	uint4 *d_input_data = d_input.resize(width*height);
+	uint4 *d_input_data = d_input.alloc(width*height);
 
 	for(int y = 0; y < height; y++) {
 		for(int x = 0; x < width; x++) {
@@ -52,16 +52,12 @@ static void shade_background_pixels(Device *device, DeviceScene *dscene, int res
 	}
 
 	/* compute on device */
-	d_output.resize(width*height);
-	memset((void*)d_output.data_pointer, 0, d_output.memory_size());
+	d_output.alloc(width*height);
+	d_output.zero_to_device();
+	d_input.copy_to_device();
 
 	device->const_copy_to("__data", &dscene->data, sizeof(dscene->data));
 
-	device->mem_alloc(d_input);
-	device->mem_copy_to(d_input);
-	device->mem_alloc(d_output);
-	device->mem_zero(d_output);
-
 	DeviceTask main_task(DeviceTask::SHADER);
 	main_task.shader_input = d_input.device_pointer;
 	main_task.shader_output = d_output.device_pointer;
@@ -78,13 +74,10 @@ static void shade_background_pixels(Device *device, DeviceScene *dscene, int res
 	foreach(DeviceTask& task, split_tasks) {
 		device->task_add(task);
 		device->task_wait();
-		device->mem_copy_from(d_output, task.shader_x, 1, task.shader_w, sizeof(float4));
+		d_output.copy_from_device(task.shader_x, 1, task.shader_w);
 	}
 
-	device->mem_free(d_input);
-	device->mem_free(d_output);
-
-	d_input.clear();
+	d_input.free();
 
 	float4 *d_output_data = reinterpret_cast<float4*>(d_output.data_pointer);
 
@@ -97,6 +90,8 @@ static void shade_background_pixels(Device *device, DeviceScene *dscene, int res
 			pixels[y*width + x].z = d_output_data[y*width + x].z;
 		}
 	}
+
+	d_output.free();
 }
 
 /* Light */
@@ -246,7 +241,7 @@ bool LightManager::object_usable_as_light(Object *object) {
 	return false;
 }
 
-void LightManager::device_update_distribution(Device *device, DeviceScene *dscene, Scene *scene, Progress& progress)
+void LightManager::device_update_distribution(Device *, DeviceScene *dscene, Scene *scene, Progress& progress)
 {
 	progress.set_status("Updating Lights", "Computing distribution");
 
@@ -292,7 +287,7 @@ void LightManager::device_update_distribution(Device *device, DeviceScene *dscen
 	VLOG(1) << "Total " << num_distribution << " of light distribution primitives.";
 
 	/* emission area */
-	float4 *distribution = dscene->light_distribution.resize(num_distribution + 1);
+	float4 *distribution = dscene->light_distribution.alloc(num_distribution + 1);
 	float totarea = 0.0f;
 
 	/* triangles */
@@ -451,7 +446,7 @@ void LightManager::device_update_distribution(Device *device, DeviceScene *dscen
 			kfilm->pass_shadow_scale *= (float)(num_lights - num_background_lights)/(float)num_lights;
 
 		/* CDF */
-		device->tex_alloc(dscene->light_distribution);
+		dscene->light_distribution.copy_to_device();
 
 		/* Portals */
 		if(num_portals > 0) {
@@ -466,7 +461,7 @@ void LightManager::device_update_distribution(Device *device, DeviceScene *dscen
 		}
 	}
 	else {
-		dscene->light_distribution.clear();
+		dscene->light_distribution.free();
 
 		kintegrator->num_distribution = 0;
 		kintegrator->num_all_lights = 0;
@@ -561,8 +556,8 @@ void LightManager::device_update_background(Device *device,
 
 	/* build row distributions and column distribution for the infinite area environment light */
 	int cdf_count = res + 1;
-	float2 *marg_cdf = dscene->light_background_marginal_cdf.resize(cdf_count);
-	float2 *cond_cdf = dscene->light_background_conditional_cdf.resize(cdf_count * cdf_count);
+	float2 *marg_cdf = dscene->light_background_marginal_cdf.alloc(cdf_count);
+	float2 *cond_cdf = dscene->light_background_conditional_cdf.alloc(cdf_count * cdf_count);
 
 	double time_start = time_dt();
 	if(res < 512) {
@@ -611,11 +606,11 @@ void LightManager::device_update_background(Device *device,
 	VLOG(2) << "Background MIS build time " << time_dt() - time_start << "\n";
 
 	/* update device */
-	device->tex_alloc(dscene->light_background_marginal_cdf);
-	device->tex_alloc(dscene->light_background_conditional_cdf);
+	dscene->light_background_marginal_cdf.copy_to_device();
+	dscene->light_background_conditional_cdf.copy_to_device();
 }
 
-void LightManager::device_update_points(Device *device,
+void LightManager::device_update_points(Device *,
                                         DeviceScene *dscene,
                                         Scene *scene)
 {
@@ -628,7 +623,7 @@ void LightManager::device_update_points(Device *device,
 		}
 	}
 
-	float4 *light_data = dscene->light_data.resize(num_lights*LIGHT_SIZE);
+	float4 *light_data = dscene->light_data.alloc(num_lights*LIGHT_SIZE);
 
 	if(num_lights == 0) {
 		VLOG(1) << "No effective light, ignoring points update.";
@@ -813,7 +808,7 @@ void LightManager::device_update_points(Device *device,
 	VLOG(1) << "Number of lights without contribution: "
 	        << num_scene_lights - light_index;
 
-	device->tex_alloc(dscene->light_data);
+	dscene->light_data.copy_to_device();
 }
 
 void LightManager::device_update(Device *device, DeviceScene *dscene, Scene *scene, Progress& progress)
@@ -846,17 +841,12 @@ void LightManager::device_update(Device *device, DeviceScene *dscene, Scene *sce
 	need_update = false;
 }
 
-void LightManager::device_free(Device *device, DeviceScene *dscene)
+void LightManager::device_free(Device *, DeviceScene *dscene)
 {
-	device->tex_free(dscene->light_distribution);
-	device->tex_free(dscene->light_data);
-	device->tex_free(dscene->light_background_marginal_cdf);
-	device->tex_free(dscene->light_background_conditional_cdf);
-
-	dscene->light_distribution.clear();
-	dscene->light_data.clear();
-	dscene->light_background_marginal_cdf.clear();
-	dscene->light_background_conditional_cdf.clear();
+	dscene->light_distribution.free();
+	dscene->light_data.free();
+	dscene->light_background_marginal_cdf.free();
+	dscene->light_background_conditional_cdf.free();
 }
 
 void LightManager::tag_update(Scene * /*scene*/)
diff --git a/intern/cycles/render/mesh.cpp b/intern/cycles/render/mesh.cpp
index 685272b80c1..75bdf71616f 100644
--- a/intern/cycles/render/mesh.cpp
+++ b/intern/cycles/render/mesh.cpp
@@ -1252,7 +1252,7 @@ void MeshManager::update_osl_attributes(Device *device, Scene *scene, vector<Att
 #endif
 }
 
-void MeshManager::update_svm_attributes(Device *device, DeviceScene *dscene, Scene *scene, vector<AttributeRequestSet>& mesh_attributes)
+void MeshManager::update_svm_attributes(Device *, DeviceScene *dscene, Scene *scene, vector<AttributeRequestSet>& mesh_attributes)
 {
 	/* for SVM, the attributes_map table is used to lookup the offset of an
 	 * attribute, based on a unique shader attribute id. */
@@ -1267,7 +1267,7 @@ void MeshManager::update_svm_attributes(Device *device, DeviceScene *dscene, Sce
 		return;
 
 	/* create attribute map */
-	uint4 *attr_map = dscene->attributes_map.resize(attr_map_stride*scene->objects.size());
+	uint4 *attr_map = dscene->attributes_map.alloc(attr_map_stride*scene->objects.size());
 	memset(attr_map, 0, dscene->attributes_map.size()*sizeof(uint));
 
 	for(size_t i = 0; i < scene->objects.size(); i++) {
@@ -1359,7 +1359,7 @@ void MeshManager::update_svm_attributes(Device *device, DeviceScene *dscene, Sce
 
 	/* copy to device */
 	dscene->data.bvh.attributes_map_stride = attr_map_stride;
-	device->tex_alloc(dscene->attributes_map);
+	dscene->attributes_map.copy_to_device();
 }
 
 static void update_attribute_element_size(Mesh *mesh,
@@ -1554,9 +1554,9 @@ void MeshManager::device_update_attributes(Device *device, DeviceScene *dscene,
 		}
 	}
 
-	dscene->attributes_float.resize(attr_float_size);
-	dscene->attributes_float3.resize(attr_float3_size);
-	dscene->attributes_uchar4.resize(attr_uchar4_size);
+	dscene->attributes_float.alloc(attr_float_size);
+	dscene->attributes_float3.alloc(attr_float3_size);
+	dscene->attributes_uchar4.alloc(attr_uchar4_size);
 
 	size_t attr_float_offset = 0;
 	size_t attr_float3_offset = 0;
@@ -1617,13 +1617,13 @@ void MeshManager::device_update_attributes(Device *device, DeviceScene *dscene,
 	progress.set_status("Updating Mesh", "Copying Attributes to device");
 
 	if(dscene->attributes_float.size()) {
-		device->tex_alloc(dscene->attributes_float);
+		dscene->attributes_float.copy_to_device();
 	}
 	if(dscene->attributes_float3.size()) {
-		device->tex_alloc(dscene->attributes_float3);
+		dscene->attributes_float3.copy_to_device();
 	}
 	if(dscene->attributes_uchar4.size()) {
-		device->tex_alloc(dscene->attributes_uchar4);
+		dscene->attributes_uchar4.copy_to_device();
 	}
 }
 
@@ -1671,7 +1671,7 @@ void MeshManager::mesh_calc_offset(Scene *scene)
 	}
 }
 
-void MeshManager::device_update_mesh(Device *device,
+void MeshManager::device_update_mesh(Device *,
                                      DeviceScene *dscene,
                                      Scene *scene,
                                      bool for_displacement,
@@ -1732,11 +1732,11 @@ void MeshManager::device_update_mesh(Device *device,
 		/* normals */
 		progress.set_status("Updating Mesh", "Computing normals");
 
-		uint *tri_shader = dscene->tri_shader.resize(tri_size);
-		float4 *vnormal = dscene->tri_vnormal.resize(vert_size);
-		uint4 *tri_vindex = dscene->tri_vindex.resize(tri_size);
-		uint *tri_patch = dscene->tri_patch.resize(tri_size);
-		float2 *tri_patch_uv = dscene->tri_patch_uv.resize(vert_size);
+		uint *tri_shader = dscene->tri_shader.alloc(tri_size);
+		float4 *vnormal = dscene->tri_vnormal.alloc(vert_size);
+		uint4 *tri_vindex = dscene->tri_vindex.alloc(tri_size);
+		uint *tri_patch = dscene->tri_patch.alloc(tri_size);
+		float2 *tri_patch_uv = dscene->tri_patch_uv.alloc(vert_size);
 
 		foreach(Mesh *mesh, scene->meshes) {
 			mesh->pack_normals(scene,
@@ -1754,32 +1754,32 @@ void MeshManager::device_update_mesh(Device *device,
 		/* vertex coordinates */
 		progress.set_status("Updating Mesh", "Copying Mesh to device");
 
-		device->tex_alloc(dscene->tri_shader);
-		device->tex_alloc(dscene->tri_vnormal);
-		device->tex_alloc(dscene->tri_vindex);
-		device->tex_alloc(dscene->tri_patch);
-		device->tex_alloc(dscene->tri_patch_uv);
+		dscene->tri_shader.copy_to_device();
+		dscene->tri_vnormal.copy_to_device();
+		dscene->tri_vindex.copy_to_device();
+		dscene->tri_patch.copy_to_device();
+		dscene->tri_patch_uv.copy_to_device();
 	}
 
 	if(curve_size != 0) {
 		progress.set_status("Updating Mesh", "Copying Strands to device");
 
-		float4 *curve_keys = dscene->curve_keys.resize(curve_key_size);
-		float4 *curves = dscene->curves.resize(curve_size);
+		float4 *curve_keys = dscene->curve_keys.alloc(curve_key_size);
+		float4 *curves = dscene->curves.alloc(curve_size);
 
 		foreach(Mesh *mesh, scene->meshes) {
 			mesh->pack_curves(scene, &curve_keys[mesh->curvekey_offset], &curves[mesh->curve_offset], mesh->curvekey_offset);
 			if(progress.get_cancel()) return;
 		}
 
-		device->tex_alloc(dscene->curve_keys);
-		device->tex_alloc(dscene->curves);
+		dscene->curve_keys.copy_to_device();
+		dscene->curves.copy_to_device();
 	}
 
 	if(patch_size != 0) {
 		progress.set_status("Updating Mesh", "Copying Patches to device");
 
-		uint *patch_data = dscene->patches.resize(patch_size);
+		uint *patch_data = dscene->patches.alloc(patch_size);
 
 		foreach(Mesh *mesh, scene->meshes) {
 			mesh->pack_patches(&patch_data[mesh->patch_offset], mesh->vert_offset, mesh->face_offset, mesh->corner_offset);
@@ -1791,11 +1791,11 @@ void MeshManager::device_update_mesh(Device *device,
 			if(progress.get_cancel()) return;
 		}
 
-		device->tex_alloc(dscene->patches);
+		dscene->patches.copy_to_device();
 	}
 
 	if(for_displacement) {
-		float4 *prim_tri_verts = dscene->prim_tri_verts.resize(tri_size * 3);
+		float4 *prim_tri_verts = dscene->prim_tri_verts.alloc(tri_size * 3);
 		foreach(Mesh *mesh, scene->meshes) {
 			for(size_t i = 0; i < mesh->num_triangles(); ++i) {
 				Mesh::Triangle t = mesh->get_triangle(i);
@@ -1805,7 +1805,7 @@ void MeshManager::device_update_mesh(Device *device,
 				prim_tri_verts[offset + 2] = float3_to_float4(mesh->verts[t.v[2]]);
 			}
 		}
-		device->tex_alloc(dscene->prim_tri_verts);
+		dscene->prim_tri_verts.copy_to_device();
 	}
 }
 
@@ -1841,43 +1841,43 @@ void MeshManager::device_update_bvh(Device *device, DeviceScene *dscene, Scene *
 
 	if(pack.nodes.size()) {
 		dscene->bvh_nodes.steal_data(pack.nodes);
-		device->tex_alloc(dscene->bvh_nodes);
+		dscene->bvh_nodes.copy_to_device();
 	}
 	if(pack.leaf_nodes.size()) {
 		dscene->bvh_leaf_nodes.steal_data(pack.leaf_nodes);
-		device->tex_alloc(dscene->bvh_leaf_nodes);
+		dscene->bvh_leaf_nodes.copy_to_device();
 	}
 	if(pack.object_node.size()) {
 		dscene->object_node.steal_data(pack.object_node);
-		device->tex_alloc(dscene->object_node);
+		dscene->object_node.copy_to_device();
 	}
 	if(pack.prim_tri_index.size()) {
 		dscene->prim_tri_index.steal_data(pack.prim_tri_index);
-		device->tex_alloc(dscene->prim_tri_index);
+		dscene->prim_tri_index.copy_to_device();
 	}
 	if(pack.prim_tri_verts.size()) {
 		dscene->prim_tri_verts.steal_data(pack.prim_tri_verts);
-		device->tex_alloc(dscene->prim_tri_verts);
+		dscene->prim_tri_verts.copy_to_device();
 	}
 	if(pack.prim_type.size()) {
 		dscene->prim_type.steal_data(pack.prim_type);
-		device->tex_alloc(dscene->prim_type);
+		dscene->prim_type.copy_to_device();
 	}
 	if(pack.prim_visibility.size()) {
 		dscene->prim_visibility.steal_data(pack.prim_visibility);
-		device->tex_alloc(dscene->prim_visibility);
+		dscene->prim_visibility.copy_to_device();
 	}
 	if(pack.prim_index.size()) {
 		dscene->prim_index.steal_data(pack.prim_index);
-		device->tex_alloc(dscene->prim_index);
+		dscene->prim_index.copy_to_device();
 	}
 	if(pack.prim_object.size()) {
 		dscene->prim_object.steal_data(pack.prim_object);
-		device->tex_alloc(dscene->prim_object);
+		dscene->prim_object.copy_to_device();
 	}
 	if(pack.prim_time.size()) {
 		dscene->prim_time.steal_data(pack.prim_time);
-		device->tex_alloc(dscene->prim_time);
+		dscene->prim_time.copy_to_device();
 	}
 
 	dscene->data.bvh.root = pack.root_index;
@@ -2142,51 +2142,28 @@ void MeshManager::device_update(Device *device, DeviceScene *dscene, Scene *scen
 
 void MeshManager::device_free(Device *device, DeviceScene *dscene)
 {
-	device->tex_free(dscene->bvh_nodes);
-	device->tex_free(dscene->bvh_leaf_nodes);
-	device->tex_free(dscene->object_node);
-	device->tex_free(dscene->prim_tri_verts);
-	device->tex_free(dscene->prim_tri_index);
-	device->tex_free(dscene->prim_type);
-	device->tex_free(dscene->prim_visibility);
-	device->tex_free(dscene->prim_index);
-	device->tex_free(dscene->prim_object);
-	device->tex_free(dscene->prim_time);
-	device->tex_free(dscene->tri_shader);
-	device->tex_free(dscene->tri_vnormal);
-	device->tex_free(dscene->tri_vindex);
-	device->tex_free(dscene->tri_patch);
-	device->tex_free(dscene->tri_patch_uv);
-	device->tex_free(dscene->curves);
-	device->tex_free(dscene->curve_keys);
-	device->tex_free(dscene->patches);
-	device->tex_free(dscene->attributes_map);
-	device->tex_free(dscene->attributes_float);
-	device->tex_free(dscene->attributes_float3);
-	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();
-	dscene->prim_type.clear();
-	dscene->prim_visibility.clear();
-	dscene->prim_index.clear();
-	dscene->prim_object.clear();
-	dscene->prim_time.clear();
-	dscene->tri_shader.clear();
-	dscene->tri_vnormal.clear();
-	dscene->tri_vindex.clear();
-	dscene->tri_patch.clear();
-	dscene->tri_patch_uv.clear();
-	dscene->curves.clear();
-	dscene->curve_keys.clear();
-	dscene->patches.clear();
-	dscene->attributes_map.clear();
-	dscene->attributes_float.clear();
-	dscene->attributes_float3.clear();
-	dscene->attributes_uchar4.clear();
+	dscene->bvh_nodes.free();
+	dscene->bvh_leaf_nodes.free();
+	dscene->object_node.free();
+	dscene->prim_tri_verts.free();
+	dscene->prim_tri_index.free();
+	dscene->prim_type.free();
+	dscene->prim_visibility.free();
+	dscene->prim_index.free();
+	dscene->prim_object.free();
+	dscene->prim_time.free();
+	dscene->tri_shader.free();
+	dscene->tri_vnormal.free();
+	dscene->tri_vindex.free();
+	dscene->tri_patch.free();
+	dscene->tri_patch_uv.free();
+	dscene->curves.free();
+	dscene->curve_keys.free();
+	dscene->patches.free();
+	dscene->attributes_map.free();
+	dscene->attributes_float.free();
+	dscene->attributes_float3.free();
+	dscene->attributes_uchar4.free();
 
 #ifdef WITH_OSL
 	OSLGlobals *og = (OSLGlobals*)device->osl_memory();
diff --git a/intern/cycles/render/mesh_displace.cpp b/intern/cycles/render/mesh_displace.cpp
index c06cf86ea9c..ab3ae40d931 100644
--- a/intern/cycles/render/mesh_displace.cpp
+++ b/intern/cycles/render/mesh_displace.cpp
@@ -65,7 +65,7 @@ bool MeshManager::displace(Device *device, DeviceScene *dscene, Scene *scene, Me
 	const size_t num_verts = mesh->verts.size();
 	vector<bool> done(num_verts, false);
 	device_vector<uint4> d_input(device, "displace_input", MEM_READ_ONLY);
-	uint4 *d_input_data = d_input.resize(num_verts);
+	uint4 *d_input_data = d_input.alloc(num_verts);
 	size_t d_input_size = 0;
 
 	size_t num_triangles = mesh->num_triangles();
@@ -116,16 +116,13 @@ bool MeshManager::displace(Device *device, DeviceScene *dscene, Scene *scene, Me
 	
 	/* run device task */
 	device_vector<float4> d_output(device, "displace_output", MEM_WRITE_ONLY);
-	d_output.resize(d_input_size);
+	d_output.alloc(d_input_size);
+	d_output.zero_to_device();
+	d_input.copy_to_device();
 
 	/* needs to be up to data for attribute access */
 	device->const_copy_to("__data", &dscene->data, sizeof(dscene->data));
 
-	device->mem_alloc(d_input);
-	device->mem_copy_to(d_input);
-	device->mem_alloc(d_output);
-	device->mem_zero(d_output);
-
 	DeviceTask task(DeviceTask::SHADER);
 	task.shader_input = d_input.device_pointer;
 	task.shader_output = d_output.device_pointer;
@@ -139,14 +136,13 @@ bool MeshManager::displace(Device *device, DeviceScene *dscene, Scene *scene, Me
 	device->task_wait();
 
 	if(progress.get_cancel()) {
-		device->mem_free(d_input);
-		device->mem_free(d_output);
+		d_input.free();
+		d_output.free();
 		return false;
 	}
 
-	device->mem_copy_from(d_output, 0, 1, d_output.size(), sizeof(float4));
-	device->mem_free(d_input);
-	device->mem_free(d_output);
+	d_output.copy_from_device(0, 1, d_output.size());
+	d_input.free();
 
 	/* read result */
 	done.clear();
@@ -183,6 +179,8 @@ bool MeshManager::displace(Device *device, DeviceScene *dscene, Scene *scene, Me
 		}
 	}
 
+	d_output.free();
+
 	/* for displacement method both, we only need to recompute the face
 	 * normals, as bump mapping in the shader will already alter the
 	 * vertex normal, so we start from the non-displaced vertex normals
diff --git a/intern/cycles/render/object.cpp b/intern/cycles/render/object.cpp
index daa872239ce..57e44861e40 100644
--- a/intern/cycles/render/object.cpp
+++ b/intern/cycles/render/object.cpp
@@ -488,9 +488,9 @@ void ObjectManager::device_update_transforms(Device *device,
 	state.queue_start_object = 0;
 
 	state.object_flag = object_flag;
-	state.objects = dscene->objects.resize(OBJECT_SIZE*scene->objects.size());
+	state.objects = dscene->objects.alloc(OBJECT_SIZE*scene->objects.size());
 	if(state.need_motion == Scene::MOTION_PASS) {
-		state.objects_vector = dscene->objects_vector.resize(OBJECT_VECTOR_SIZE*scene->objects.size());
+		state.objects_vector = dscene->objects_vector.alloc(OBJECT_VECTOR_SIZE*scene->objects.size());
 	}
 	else {
 		state.objects_vector = NULL;
@@ -534,9 +534,9 @@ void ObjectManager::device_update_transforms(Device *device,
 		}
 	}
 
-	device->tex_alloc(dscene->objects);
+	dscene->objects.copy_to_device();
 	if(state.need_motion == Scene::MOTION_PASS) {
-		device->tex_alloc(dscene->objects_vector);
+		dscene->objects_vector.copy_to_device();
 	}
 
 	dscene->data.bvh.have_motion = state.have_motion;
@@ -557,7 +557,7 @@ void ObjectManager::device_update(Device *device, DeviceScene *dscene, Scene *sc
 		return;
 
 	/* object info flag */
-	uint *object_flag = dscene->object_flag.resize(scene->objects.size());
+	uint *object_flag = dscene->object_flag.alloc(scene->objects.size());
 
 	/* set object transform matrices, before applying static transforms */
 	progress.set_status("Updating Objects", "Copying Transformations to device");
@@ -573,7 +573,7 @@ void ObjectManager::device_update(Device *device, DeviceScene *dscene, Scene *sc
 	}
 }
 
-void ObjectManager::device_update_flags(Device *device,
+void ObjectManager::device_update_flags(Device *,
                                         DeviceScene *dscene,
                                         Scene *scene,
                                         Progress& /*progress*/,
@@ -638,10 +638,10 @@ void ObjectManager::device_update_flags(Device *device,
 	}
 
 	/* allocate object flag */
-	device->tex_alloc(dscene->object_flag);
+	dscene->object_flag.copy_to_device();
 }
 
-void ObjectManager::device_update_patch_map_offsets(Device *device, DeviceScene *dscene, Scene *scene)
+void ObjectManager::device_update_patch_map_offsets(Device *, DeviceScene *dscene, Scene *scene)
 {
 	if(scene->objects.size() == 0) {
 		return;
@@ -671,21 +671,15 @@ void ObjectManager::device_update_patch_map_offsets(Device *device, DeviceScene
 	}
 
 	if(update) {
-		device->tex_free(dscene->objects);
-		device->tex_alloc(dscene->objects);
+		dscene->objects.copy_to_device();
 	}
 }
 
-void ObjectManager::device_free(Device *device, DeviceScene *dscene)
+void ObjectManager::device_free(Device *, DeviceScene *dscene)
 {
-	device->tex_free(dscene->objects);
-	dscene->objects.clear();
-
-	device->tex_free(dscene->objects_vector);
-	dscene->objects_vector.clear();
-
-	device->tex_free(dscene->object_flag);
-	dscene->object_flag.clear();
+	dscene->objects.free();
+	dscene->objects_vector.free();
+	dscene->object_flag.free();
 }
 
 void ObjectManager::apply_static_transforms(DeviceScene *dscene, Scene *scene, uint *object_flag, Progress& progress)
diff --git a/intern/cycles/render/particles.cpp b/intern/cycles/render/particles.cpp
index a84ca51f274..06ff45b09bd 100644
--- a/intern/cycles/render/particles.cpp
+++ b/intern/cycles/render/particles.cpp
@@ -52,7 +52,7 @@ ParticleSystemManager::~ParticleSystemManager()
 {
 }
 
-void ParticleSystemManager::device_update_particles(Device *device, DeviceScene *dscene, Scene *scene, Progress& progress)
+void ParticleSystemManager::device_update_particles(Device *, DeviceScene *dscene, Scene *scene, Progress& progress)
 {
 	/* count particles.
 	 * adds one dummy particle at the beginning to avoid invalid lookups,
@@ -61,7 +61,7 @@ void ParticleSystemManager::device_update_particles(Device *device, DeviceScene
 	for(size_t j = 0; j < scene->particle_systems.size(); j++)
 		num_particles += scene->particle_systems[j]->particles.size();
 	
-	float4 *particles = dscene->particles.resize(PARTICLE_SIZE*num_particles);
+	float4 *particles = dscene->particles.alloc(PARTICLE_SIZE*num_particles);
 	
 	/* dummy particle */
 	particles[0] = make_float4(0.0f, 0.0f, 0.0f, 0.0f);
@@ -91,7 +91,7 @@ void ParticleSystemManager::device_update_particles(Device *device, DeviceScene
 		}
 	}
 	
-	device->tex_alloc(dscene->particles);
+	dscene->particles.copy_to_device();
 }
 
 void ParticleSystemManager::device_update(Device *device, DeviceScene *dscene, Scene *scene, Progress& progress)
@@ -112,10 +112,9 @@ void ParticleSystemManager::device_update(Device *device, DeviceScene *dscene, S
 	need_update = false;
 }
 
-void ParticleSystemManager::device_free(Device *device, DeviceScene *dscene)
+void ParticleSystemManager::device_free(Device *, DeviceScene *dscene)
 {
-	device->tex_free(dscene->particles);
-	dscene->particles.clear();
+	dscene->particles.free();
 }
 
 void ParticleSystemManager::tag_update(Scene * /*scene*/)
diff --git a/intern/cycles/render/scene.cpp b/intern/cycles/render/scene.cpp
index e362a35471d..260a325206c 100644
--- a/intern/cycles/render/scene.cpp
+++ b/intern/cycles/render/scene.cpp
@@ -41,40 +41,40 @@
 CCL_NAMESPACE_BEGIN
 
 DeviceScene::DeviceScene(Device *device)
-: bvh_nodes(device, "__bvh_nodes"),
-  bvh_leaf_nodes(device, "__bvh_leaf_nodes"),
-  object_node(device, "__object_node"),
-  prim_tri_index(device, "__prim_tri_index"),
-  prim_tri_verts(device, "__prim_tri_verts"),
-  prim_type(device, "__prim_type"),
-  prim_visibility(device, "__prim_visibility"),
-  prim_index(device, "__prim_index"),
-  prim_object(device, "__prim_object"),
-  prim_time(device, "__prim_time"),
-  tri_shader(device, "__tri_shader"),
-  tri_vnormal(device, "__tri_vnormal"),
-  tri_vindex(device, "__tri_vindex"),
-  tri_patch(device, "__tri_patch"),
-  tri_patch_uv(device, "__tri_patch_uv"),
-  curves(device, "__curves"),
-  curve_keys(device, "__curve_keys"),
-  patches(device, "__patches"),
-  objects(device, "__objects"),
-  objects_vector(device, "__objects_vector"),
-  attributes_map(device, "__attributes_map"),
-  attributes_float(device, "__attributes_float"),
-  attributes_float3(device, "__attributes_float3"),
-  attributes_uchar4(device, "__attributes_uchar4"),
-  light_distribution(device, "__light_distribution"),
-  light_data(device, "__light_data"),
-  light_background_marginal_cdf(device, "__light_background_marginal_cdf"),
-  light_background_conditional_cdf(device, "__light_background_conditional_cdf"),
-  particles(device, "__particles"),
-  svm_nodes(device, "__svm_nodes"),
-  shader_flag(device, "__shader_flag"),
-  object_flag(device, "__object_flag"),
-  lookup_table(device, "__lookup_table"),
-  sobol_directions(device, "__sobol_directions")
+: bvh_nodes(device, "__bvh_nodes", MEM_TEXTURE),
+  bvh_leaf_nodes(device, "__bvh_leaf_nodes", MEM_TEXTURE),
+  object_node(device, "__object_node", MEM_TEXTURE),
+  prim_tri_index(device, "__prim_tri_index", MEM_TEXTURE),
+  prim_tri_verts(device, "__prim_tri_verts", MEM_TEXTURE),
+  prim_type(device, "__prim_type", MEM_TEXTURE),
+  prim_visibility(device, "__prim_visibility", MEM_TEXTURE),
+  prim_index(device, "__prim_index", MEM_TEXTURE),
+  prim_object(device, "__prim_object", MEM_TEXTURE),
+  prim_time(device, "__prim_time", MEM_TEXTURE),
+  tri_shader(device, "__tri_shader", MEM_TEXTURE),
+  tri_vnormal(device, "__tri_vnormal", MEM_TEXTURE),
+  tri_vindex(device, "__tri_vindex", MEM_TEXTURE),
+  tri_patch(device, "__tri_patch", MEM_TEXTURE),
+  tri_patch_uv(device, "__tri_patch_uv", MEM_TEXTURE),
+  curves(device, "__curves", MEM_TEXTURE),
+  curve_keys(device, "__curve_keys", MEM_TEXTURE),
+  patches(device, "__patches", MEM_TEXTURE),
+  objects(device, "__objects", MEM_TEXTURE),
+  objects_vector(device, "__objects_vector", MEM_TEXTURE),
+  attributes_map(device, "__attributes_map", MEM_TEXTURE),
+  attributes_float(device, "__attributes_float", MEM_TEXTURE),
+  attributes_float3(device, "__attributes_float3", MEM_TEXTURE),
+  attributes_uchar4(device, "__attributes_uchar4", MEM_TEXTURE),
+  light_distribution(device, "__light_distribution", MEM_TEXTURE),
+  light_data(device, "__light_data", MEM_TEXTURE),
+  light_background_marginal_cdf(device, "__light_background_marginal_cdf", MEM_TEXTURE),
+  light_background_conditional_cdf(device, "__light_background_conditional_cdf", MEM_TEXTURE),
+  particles(device, "__particles", MEM_TEXTURE),
+  svm_nodes(device, "__svm_nodes", MEM_TEXTURE),
+  shader_flag(device, "__shader_flag", MEM_TEXTURE),
+  object_flag(device, "__object_flag", MEM_TEXTURE),
+  lookup_table(device, "__lookup_table", MEM_TEXTURE),
+  sobol_directions(device, "__sobol_directions", MEM_TEXTURE)
 {
 	memset(&data, 0, sizeof(data));
 }
diff --git a/intern/cycles/render/session.cpp b/intern/cycles/render/session.cpp
index 4642dcfa9a1..74cfd02e1a4 100644
--- a/intern/cycles/render/session.cpp
+++ b/intern/cycles/render/session.cpp
@@ -106,11 +106,11 @@ Session::~Session()
 		delete display;
 
 		display = new DisplayBuffer(device, false);
-		display->reset(device, buffers->params);
+		display->reset(buffers->params);
 		tonemap(params.samples);
 
 		progress.set_status("Writing Image", params.output_path);
-		display->write(device, params.output_path);
+		display->write(params.output_path);
 	}
 
 	/* clean up */
@@ -399,7 +399,7 @@ bool Session::acquire_tile(Device *tile_device, RenderTile& rtile)
 
 		/* allocate buffers */
 		tile->buffers = new RenderBuffers(tile_device);
-		tile->buffers->reset(tile_device, buffer_params);
+		tile->buffers->reset(buffer_params);
 	}
 
 	tile->buffers->params.get_offset_stride(rtile.offset, rtile.stride);
@@ -756,9 +756,9 @@ void Session::reset_(BufferParams& buffer_params, int samples)
 {
 	if(buffers && buffer_params.modified(tile_manager.params)) {
 		gpu_draw_ready = false;
-		buffers->reset(device, buffer_params);
+		buffers->reset(buffer_params);
 		if(display) {
-			display->reset(device, buffer_params);
+			display->reset(buffer_params);
 		}
 	}
 
@@ -923,7 +923,7 @@ void Session::render()
 {
 	/* Clear buffers. */
 	if(buffers && tile_manager.state.sample == tile_manager.range_start_sample) {
-		buffers->zero(device);
+		buffers->zero();
 	}
 
 	/* Add path trace task. */
diff --git a/intern/cycles/render/shader.cpp b/intern/cycles/render/shader.cpp
index a77df55e520..70f6d5bab47 100644
--- a/intern/cycles/render/shader.cpp
+++ b/intern/cycles/render/shader.cpp
@@ -416,14 +416,13 @@ void ShaderManager::device_update_common(Device *device,
                                          Scene *scene,
                                          Progress& /*progress*/)
 {
-	device->tex_free(dscene->shader_flag);
-	dscene->shader_flag.clear();
+	dscene->shader_flag.free();
 
 	if(scene->shaders.size() == 0)
 		return;
 
 	uint shader_flag_size = scene->shaders.size()*SHADER_SIZE;
-	uint *shader_flag = dscene->shader_flag.resize(shader_flag_size);
+	uint *shader_flag = dscene->shader_flag.alloc(shader_flag_size);
 	uint i = 0;
 	bool has_volumes = false;
 	bool has_transparent_shadow = false;
@@ -479,7 +478,7 @@ void ShaderManager::device_update_common(Device *device,
 		has_transparent_shadow |= (flag & SD_HAS_TRANSPARENT_SHADOW) != 0;
 	}
 
-	device->tex_alloc(dscene->shader_flag);
+	dscene->shader_flag.copy_to_device();
 
 	/* lookup tables */
 	KernelTables *ktables = &dscene->data.tables;
@@ -504,12 +503,11 @@ void ShaderManager::device_update_common(Device *device,
 	kintegrator->transparent_shadows = has_transparent_shadow;
 }
 
-void ShaderManager::device_free_common(Device *device, DeviceScene *dscene, Scene *scene)
+void ShaderManager::device_free_common(Device *, DeviceScene *dscene, Scene *scene)
 {
 	scene->lookup_tables->remove_table(&beckmann_table_offset);
 
-	device->tex_free(dscene->shader_flag);
-	dscene->shader_flag.clear();
+	dscene->shader_flag.free();
 }
 
 void ShaderManager::add_default(Scene *scene)
diff --git a/intern/cycles/render/svm.cpp b/intern/cycles/render/svm.cpp
index cf0dc97ef3f..db53e366d1e 100644
--- a/intern/cycles/render/svm.cpp
+++ b/intern/cycles/render/svm.cpp
@@ -130,7 +130,7 @@ void SVMShaderManager::device_update(Device *device, DeviceScene *dscene, Scene
 	}
 
 	dscene->svm_nodes.steal_data(svm_nodes);
-	device->tex_alloc(dscene->svm_nodes);
+	dscene->svm_nodes.copy_to_device();
 
 	for(i = 0; i < scene->shaders.size(); i++) {
 		Shader *shader = scene->shaders[i];
@@ -150,8 +150,7 @@ void SVMShaderManager::device_free(Device *device, DeviceScene *dscene, Scene *s
 {
 	device_free_common(device, dscene, scene);
 
-	device->tex_free(dscene->svm_nodes);
-	dscene->svm_nodes.clear();
+	dscene->svm_nodes.free();
 }
 
 /* Graph Compiler */
diff --git a/intern/cycles/render/tables.cpp b/intern/cycles/render/tables.cpp
index 9d04778abc6..5cda977b7f1 100644
--- a/intern/cycles/render/tables.cpp
+++ b/intern/cycles/render/tables.cpp
@@ -35,25 +35,22 @@ LookupTables::~LookupTables()
 	assert(lookup_tables.size() == 0);
 }
 
-void LookupTables::device_update(Device *device, DeviceScene *dscene)
+void LookupTables::device_update(Device *, DeviceScene *dscene)
 {
 	if(!need_update)
 		return;
 
 	VLOG(1) << "Total " << lookup_tables.size() << " lookup tables.";
 
-	device->tex_free(dscene->lookup_table);
-
 	if(lookup_tables.size() > 0)
-		device->tex_alloc(dscene->lookup_table);
+		dscene->lookup_table.copy_to_device();
 
 	need_update = false;
 }
 
-void LookupTables::device_free(Device *device, DeviceScene *dscene)
+void LookupTables::device_free(Device *, DeviceScene *dscene)
 {
-	device->tex_free(dscene->lookup_table);
-	dscene->lookup_table.clear();
+	dscene->lookup_table.free();
 }
 
 static size_t round_up_to_multiple(size_t size, size_t chunk)
diff --git a/intern/cycles/util/util_vector.h b/intern/cycles/util/util_vector.h
index 9e74505b14a..ca6b56c9c7e 100644
--- a/intern/cycles/util/util_vector.h
+++ b/intern/cycles/util/util_vector.h
@@ -177,6 +177,14 @@ public:
 		}
 	}
 
+	T *steal_pointer()
+	{
+		T *ptr = data_;
+		data_ = NULL;
+		clear();
+		return ptr;
+	}
+
 	T* resize(size_t newsize)
 	{
 		if(newsize == 0) {