Clean up. Change VDB iteration method in Cycles to leaf iterators rather than accessors.

10 files changed, 544 insertions, 398 deletions

diff --git a/intern/cycles/device/device_cpu.cpp b/intern/cycles/device/device_cpu.cpp
index f29b54ff237..0b8e95c03c9 100644
--- a/intern/cycles/device/device_cpu.cpp
+++ b/intern/cycles/device/device_cpu.cpp
@@ -377,9 +377,11 @@ public:
 	void tex_alloc(device_memory& mem)
 	{
 		size_t total_memory = mem.memory_size();
-		device_memory *grid_info = mem.grid_info;
-		if(grid_info) {
-			total_memory += grid_info->memory_size();
+		device_memory *sparse_mem = NULL;
+
+		if(mem.grid_info && mem.grid_type == IMAGE_GRID_TYPE_SPARSE) {
+			sparse_mem = (device_memory*)mem.grid_info;
+			total_memory += sparse_mem->memory_size();
 		}
 
 		VLOG(1) << "Texture allocate: " << mem.name << ", "
@@ -420,16 +422,12 @@ public:
 			info.height = mem.real_height;
 			info.depth = mem.real_depth;
 
-			/* For OpenVDB textures, the kernel will retrieve the accessor from
-			 * util, but there must be some value stored in data or the texture
-			 * will not be used. As a stopgap, the accessor pointer will just
-			 * be stored in both data and util. */
 			switch(mem.grid_type) {
 				case IMAGE_GRID_TYPE_OPENVDB:
-					info.util = info.data;
+					info.util = (uint64_t)mem.grid_info;
 					break;
 				case IMAGE_GRID_TYPE_SPARSE:
-					info.util = (uint64_t)grid_info->host_pointer;
+					info.util = (uint64_t)sparse_mem->host_pointer;
 					info.tiled_width = get_tile_res(info.width);
 					info.tiled_height = get_tile_res(info.height);
 					info.even_width = info.width - (info.width % TILE_SIZE);
@@ -449,18 +447,19 @@ public:
 		mem.device_size = mem.memory_size();
 		stats.mem_alloc(mem.device_size);
 
-		if(grid_info) {
-			grid_info->device_pointer = (device_ptr)grid_info->host_pointer;
-			grid_info->device_size = grid_info->memory_size();
-			stats.mem_alloc(grid_info->device_size);
+		if(sparse_mem) {
+			sparse_mem->device_pointer = (device_ptr)sparse_mem->host_pointer;
+			sparse_mem->device_size = sparse_mem->memory_size();
+			stats.mem_alloc(sparse_mem->device_size);
 		}
 	}
 
 	void tex_free(device_memory& mem)
 	{
 		if(mem.device_pointer) {
-			if(mem.grid_info) {
-				tex_free(*mem.grid_info);
+			if(mem.grid_info && mem.grid_type == IMAGE_GRID_TYPE_SPARSE) {
+				device_memory *grid_info = (device_memory*)mem.grid_info;
+				tex_free(*grid_info);
 			}
 			mem.device_pointer = 0;
 			stats.mem_free(mem.device_size);
diff --git a/intern/cycles/device/device_memory.cpp b/intern/cycles/device/device_memory.cpp
index c6248fcf88b..6e5ded19699 100644
--- a/intern/cycles/device/device_memory.cpp
+++ b/intern/cycles/device/device_memory.cpp
@@ -36,7 +36,8 @@ device_memory::device_memory(Device *device, const char *name, MemoryType type)
   device(device),
   device_pointer(0),
   host_pointer(0),
-  shared_pointer(0)
+  shared_pointer(0),
+  grid_info(0)
 {
 }
 
diff --git a/intern/cycles/device/device_memory.h b/intern/cycles/device/device_memory.h
index e21636009ee..5d4fb130ca2 100644
--- a/intern/cycles/device/device_memory.h
+++ b/intern/cycles/device/device_memory.h
@@ -207,7 +207,7 @@ public:
 	device_ptr device_pointer;
 	void *host_pointer;
 	void *shared_pointer;
-	device_memory *grid_info = NULL;
+	void *grid_info;
 
 	virtual ~device_memory();
 
diff --git a/intern/cycles/device/device_memory_openvdb.h b/intern/cycles/device/device_memory_openvdb.h
index eb9fbe4912b..ee655318b61 100644
--- a/intern/cycles/device/device_memory_openvdb.h
+++ b/intern/cycles/device/device_memory_openvdb.h
@@ -41,7 +41,8 @@ public:
 
 		assert(data_elements > 0);
 
-		host_pointer = static_cast<void*>(&vdb_acc);
+		host_pointer = static_cast<void*>(&vdb_grid);
+		grid_info = static_cast<void*>(&vdb_acc);
 
 		data_width = real_width = resolution.x;
 		data_height = real_height = resolution.y;
diff --git a/intern/cycles/render/image.cpp b/intern/cycles/render/image.cpp
index 13892db1364..8c79d245614 100644
--- a/intern/cycles/render/image.cpp
+++ b/intern/cycles/render/image.cpp
@@ -474,6 +474,33 @@ void ImageManager::tag_reload_image(const string& filename,
 	}
 }
 
+bool ImageManager::allocate_sparse_index(Device *device,
+                                         device_memory *tex_img,
+                                         vector<int> *sparse_index,
+                                         string mem_name)
+{
+	mem_name += "_index";
+	device_vector<int> *tex_index =
+	        new device_vector<int>(device, mem_name.c_str(), MEM_TEXTURE);
+
+	int *ti;
+	{
+		thread_scoped_lock device_lock(device_mutex);
+		ti = (int*)tex_index->alloc(sparse_index->size());
+	}
+
+	if(ti == NULL) {
+		return false;
+	}
+
+	memcpy(ti, &(*sparse_index)[0], sparse_index->size() * sizeof(int));
+
+	tex_img->grid_info = static_cast<void*>(tex_index);
+	tex_img->grid_type = IMAGE_GRID_TYPE_SPARSE;
+
+	return true;
+}
+
 bool ImageManager::file_load_image_generic(Image *img,
                                            ImageInput **in,
                                            int &width,
@@ -565,8 +592,7 @@ bool ImageManager::file_load_image_generic(Image *img,
 	return true;
 }
 
-template<typename StorageType,
-         typename DeviceType>
+template<typename DeviceType>
 void ImageManager::file_load_failed(Image *img,
                                     ImageDataType type,
 			                        device_vector<DeviceType> *tex_img)
@@ -577,36 +603,54 @@ void ImageManager::file_load_failed(Image *img,
 
 	/* On failure to load, we set a 1x1 pixels pink image. */
 	thread_scoped_lock device_lock(device_mutex);
-	StorageType *pixels = (StorageType*)tex_img->alloc(1, 1);
+	DeviceType *device_pixels = tex_img->alloc(1, 1);
 
 	switch(type) {
 		case IMAGE_DATA_TYPE_FLOAT4:
-			pixels[0] = TEX_IMAGE_MISSING_R;
-			pixels[1] = TEX_IMAGE_MISSING_G;
-			pixels[2] = TEX_IMAGE_MISSING_B;
-			pixels[3] = TEX_IMAGE_MISSING_A;
+		{
+			float4 *pixels = (float4*)device_pixels;
+			pixels[0].x = TEX_IMAGE_MISSING_R;
+			pixels[0].y = TEX_IMAGE_MISSING_G;
+			pixels[0].z = TEX_IMAGE_MISSING_B;
+			pixels[0].w = TEX_IMAGE_MISSING_A;
 			break;
+		}
 		case IMAGE_DATA_TYPE_FLOAT:
+		{
+			float *pixels = (float*)device_pixels;
 			pixels[0] = TEX_IMAGE_MISSING_R;
 			break;
+		}
 		case IMAGE_DATA_TYPE_BYTE4:
-			pixels[0] = (TEX_IMAGE_MISSING_R * 255);
-			pixels[1] = (TEX_IMAGE_MISSING_G * 255);
-			pixels[2] = (TEX_IMAGE_MISSING_B * 255);
-			pixels[3] = (TEX_IMAGE_MISSING_A * 255);
+		{
+			uchar4 *pixels = (uchar4*)device_pixels;
+			pixels[0].x = (TEX_IMAGE_MISSING_R * 255);
+			pixels[0].y = (TEX_IMAGE_MISSING_G * 255);
+			pixels[0].z = (TEX_IMAGE_MISSING_B * 255);
+			pixels[0].w = (TEX_IMAGE_MISSING_A * 255);
 			break;
+		}
 		case IMAGE_DATA_TYPE_BYTE:
+		{
+			uchar *pixels = (uchar*)device_pixels;
 			pixels[0] = (TEX_IMAGE_MISSING_R * 255);
 			break;
+		}
 		case IMAGE_DATA_TYPE_HALF4:
-			pixels[0] = TEX_IMAGE_MISSING_R;
-			pixels[1] = TEX_IMAGE_MISSING_G;
-			pixels[2] = TEX_IMAGE_MISSING_B;
-			pixels[3] = TEX_IMAGE_MISSING_A;
+		{
+			half4 *pixels = (half4*)device_pixels;
+			pixels[0].x = TEX_IMAGE_MISSING_R;
+			pixels[0].y = TEX_IMAGE_MISSING_G;
+			pixels[0].z = TEX_IMAGE_MISSING_B;
+			pixels[0].w = TEX_IMAGE_MISSING_A;
 			break;
+		}
 		case IMAGE_DATA_TYPE_HALF:
+		{
+			half *pixels = (half*)device_pixels;
 			pixels[0] = TEX_IMAGE_MISSING_R;
 			break;
+		}
 		default:
 			assert(0);
 	}
@@ -615,9 +659,115 @@ void ImageManager::file_load_failed(Image *img,
 	img->mem = tex_img;
 	img->mem->interpolation = img->interpolation;
 	img->mem->extension = img->extension;
+	img->mem->grid_type = IMAGE_GRID_TYPE_DEFAULT;
+
 	tex_img->copy_to_device();
 }
 
+#ifdef WITH_OPENVDB
+template<typename DeviceType>
+void ImageManager::file_load_extern_vdb(Device *device,
+                                        Image *img,
+                                        ImageDataType type)
+{
+	VLOG(1) << "Loading external VDB " << img->filename
+	        << ", Grid: " << img->grid_name;
+
+	device_vector<DeviceType> *tex_img =
+	        new device_vector<DeviceType>(device,
+	                                      img->mem_name.c_str(),
+	                                      MEM_TEXTURE);
+
+	/* Retrieve metadata. */
+	int width, height, depth, components;
+	if(!file_load_image_generic(img, NULL, width, height, depth, components)) {
+		file_load_failed<DeviceType>(img, type, tex_img);
+		return;
+	}
+
+	if(device->info.type == DEVICE_CPU && 0) {
+		/* Load pointer to OpenVDB grid into texture. */
+		device_memory *tex_vdb = NULL;
+		{
+			thread_scoped_lock device_lock(device_mutex);
+			tex_vdb = openvdb_load_device_extern(device,
+												 img->filename,
+												 img->grid_name,
+												 img->mem_name,
+												 img->interpolation,
+												 img->extension,
+												 components > 1);
+
+		}
+
+		if(tex_vdb) {
+			img->mem = tex_vdb;
+			delete tex_img;
+			tex_img = NULL;
+		}
+		else {
+			file_load_failed<DeviceType>(img, type, tex_img);
+		}
+	}
+	else {
+		/* Load VDB as device_vector. */
+		int sparse_size = -1;
+		vector<int> sparse_index;
+		openvdb_load_preprocess(img->filename, img->grid_name, components,
+		                        img->isovalue, &sparse_index, sparse_size);
+
+		/* Allocate space for image. */
+		float *pixels;
+		{
+			thread_scoped_lock device_lock(device_mutex);
+			if(sparse_size > -1) {
+				pixels = (float*)tex_img->alloc(sparse_size);
+			}
+			else {
+				pixels = (float*)tex_img->alloc(width, height, depth);
+			}
+		}
+
+		if(!pixels) {
+			/* Could be that we've run out of memory. */
+			file_load_failed<DeviceType>(img, type, tex_img);
+			return;
+		}
+
+		/* Load image. */
+		openvdb_load_image(img->filename, img->grid_name, components, pixels, &sparse_index);
+
+		/* Allocate space for sparse_index if it exists. */
+		if(sparse_size > -1) {
+			tex_img->grid_type = IMAGE_GRID_TYPE_SPARSE;
+
+			if(!allocate_sparse_index(device, (device_memory*)tex_img,
+			                          &sparse_index, img->mem_name))
+			{
+				/* Could be that we've run out of memory. */
+				file_load_failed<DeviceType>(img, type, tex_img);
+				return;
+			}
+		}
+		else {
+			tex_img->grid_type = IMAGE_GRID_TYPE_DEFAULT;
+		}
+
+		/* Set metadata and copy. */
+		tex_img->real_width = width;
+		tex_img->real_height = height;
+		tex_img->real_depth = depth;
+		tex_img->interpolation = img->interpolation;
+		tex_img->extension = img->extension;
+
+		img->mem = tex_img;
+
+		thread_scoped_lock device_lock(device_mutex);
+		tex_img->copy_to_device();
+	}
+}
+#endif
+
 template<TypeDesc::BASETYPE FileFormat,
          typename StorageType,
          typename DeviceType>
@@ -630,7 +780,10 @@ void ImageManager::file_load_image(Device *device,
 	        new device_vector<DeviceType>(device,
 	                                      img->mem_name.c_str(),
 	                                      MEM_TEXTURE);
+
 	tex_img->grid_type = IMAGE_GRID_TYPE_DEFAULT;
+	tex_img->interpolation = img->interpolation;
+	tex_img->extension = img->extension;
 
 	/* Try to retrieve an ImageInput for reading the image.
 	 * Otherwise, retrieve metadata. */
@@ -638,7 +791,7 @@ void ImageManager::file_load_image(Device *device,
 	int width, height, depth, components;
 	if(!file_load_image_generic(img, &in, width, height, depth, components)) {
 		/* Could not retrieve image. */
-		file_load_failed<StorageType, DeviceType>(img, type, tex_img);
+		file_load_failed<DeviceType>(img, type, tex_img);
 		return;
 	}
 
@@ -646,46 +799,9 @@ void ImageManager::file_load_image(Device *device,
 	size_t num_pixels = ((size_t)width) * height * depth;
 	if(max_size == 0) {
 		/* Don't bother with invalid images. */
-		file_load_failed<StorageType, DeviceType>(img, type, tex_img);
-		return;
-	}
-
-	const bool is_rgba = (type == IMAGE_DATA_TYPE_FLOAT4 ||
-	                      type == IMAGE_DATA_TYPE_HALF4 ||
-	                      type == IMAGE_DATA_TYPE_BYTE4);
-	const bool is_extern_vdb = string_endswith(img->filename, ".vdb");
-
-#ifdef WITH_OPENVDB
-	/* Load pointer to OpenVDB grid into texture if using CPU. */
-	if(is_extern_vdb && device->info.type == DEVICE_CPU && 0) {
-		VLOG(1) << "Loading external VDB " << img->filename
-		        << ", Grid: " << img->grid_name;
-
-		device_memory *tex_vdb = NULL;
-		{
-			thread_scoped_lock device_lock(device_mutex);
-			tex_vdb = openvdb_load_device_extern(device,
-			                                     img->filename,
-			                                     img->grid_name,
-			                                     img->mem_name,
-			                                     img->interpolation,
-			                                     img->extension,
-			                                     is_rgba);
-		}
-
-		if(tex_vdb) {
-			img->mem = tex_vdb;
-			delete tex_img;
-			tex_img = NULL;
-		}
-		else {
-			file_load_failed<StorageType, DeviceType>(img, type, tex_img);
-		}
-		/* For now, if we directly load OpenVDB grids, do not do any image
-		 * post-processing. */
+		file_load_failed<DeviceType>(img, type, tex_img);
 		return;
 	}
-#endif
 
 	/* Allocate storage for the image. */
 	vector<StorageType> pixels_storage;
@@ -701,7 +817,7 @@ void ImageManager::file_load_image(Device *device,
 
 	if(pixels == NULL) {
 		/* Could be that we've run out of memory. */
-		file_load_failed<StorageType, DeviceType>(img, type, tex_img);
+		file_load_failed<DeviceType>(img, type, tex_img);
 		return;
 	}
 
@@ -737,13 +853,6 @@ void ImageManager::file_load_image(Device *device,
 		in->close();
 		delete in;
 	}
-#ifdef WITH_OPENVDB
-	else if(is_extern_vdb) {
-		VLOG(1) << "Loading external VDB " << img->filename
-		        << ", Grid: " << img->grid_name;
-		openvdb_load_dense(img->filename, img->grid_name, (float*)pixels, components);
-	}
-#endif
 	else {
 		if(FileFormat == TypeDesc::FLOAT) {
 			builtin_image_float_pixels_cb(img->filename,
@@ -771,6 +880,9 @@ void ImageManager::file_load_image(Device *device,
 	 */
 	const StorageType alpha_one = (FileFormat == TypeDesc::UINT8)? 255 : 1;
 	const bool cmyk = (in ? strcmp(in->format_name(), "jpeg") == 0 && components == 4 : false);
+	const bool is_rgba = (type == IMAGE_DATA_TYPE_FLOAT4 ||
+	                      type == IMAGE_DATA_TYPE_HALF4 ||
+	                      type == IMAGE_DATA_TYPE_BYTE4);
 
 	if(is_rgba) {
 		if(cmyk) {
@@ -874,28 +986,19 @@ void ImageManager::file_load_image(Device *device,
 
 	/* Compress image if needed. */
 	int num_pixels_real = -1;
-	device_vector<int> *tex_info = NULL;
-
 	if(img->is_volume && device->info.type != DEVICE_CUDA) {
 		vector<StorageType> sparse_pixels;
-		vector<int> grid_info;
+		vector<int> sparse_index;
 
 		if(create_sparse_grid<StorageType>(pixels, width, height, depth,
 		                                   components, img->filename,
-	                                       img->isovalue, &sparse_pixels, &grid_info))
+		                                   img->isovalue, &sparse_pixels,
+		                                   &sparse_index))
 		{
 			pixels = &sparse_pixels[0];
 			num_pixels_real = sparse_pixels.size() / components;
-
-			tex_info = new device_vector<int>(device,
-			                                  (img->mem_name + "_info").c_str(),
-			                                  MEM_TEXTURE);
-			int *texture_info;
-			{
-				thread_scoped_lock device_lock(device_mutex);
-				texture_info = (int*)tex_info->alloc(grid_info.size());
-			}
-			memcpy(texture_info, &grid_info[0], grid_info.size() * sizeof(int));
+			allocate_sparse_index(device, (device_memory*)tex_img,
+			                      &sparse_index, img->mem_name);
 		}
 	}
 
@@ -919,13 +1022,6 @@ void ImageManager::file_load_image(Device *device,
 	tex_img->real_width = width;
 	tex_img->real_height = height;
 	tex_img->real_depth = depth;
-	tex_img->interpolation = img->interpolation;
-	tex_img->extension = img->extension;
-
-	if(tex_info != NULL) {
-		tex_img->grid_info = tex_info;
-		tex_img->grid_type = IMAGE_GRID_TYPE_SPARSE;
-	}
 
 	img->mem = tex_img;
 
@@ -957,8 +1053,9 @@ void ImageManager::device_load_image(Device *device,
 	/* Free previous texture(s) in slot. */
 	if(img->mem) {
 		thread_scoped_lock device_lock(device_mutex);
-		if(img->mem->grid_info) {
-			delete img->mem->grid_info;
+		if(img->mem->grid_info && img->mem->grid_type == IMAGE_GRID_TYPE_SPARSE) {
+			device_memory *info = (device_memory*)img->mem->grid_info;
+			delete info;
 			img->mem->grid_info = NULL;
 		}
 		delete img->mem;
@@ -967,13 +1064,24 @@ void ImageManager::device_load_image(Device *device,
 
 	/* Create new texture. */
 	const int texture_limit = scene->params.texture_limit;
+	const bool is_extern_vdb = string_endswith(img->filename, ".vdb");
 
 	switch(type) {
 		case IMAGE_DATA_TYPE_FLOAT4:
-			file_load_image<TypeDesc::FLOAT, float, float4>(device, img, type, texture_limit);
+#ifdef WITH_OPENVDB
+			if(is_extern_vdb)
+				file_load_extern_vdb<float4>(device, img, type);
+			else
+#endif
+				file_load_image<TypeDesc::FLOAT, float, float4>(device, img, type, texture_limit);
 			break;
 		case IMAGE_DATA_TYPE_FLOAT:
-			file_load_image<TypeDesc::FLOAT, float, float>(device, img, type, texture_limit);
+#ifdef WITH_OPENVDB
+			if(is_extern_vdb)
+				file_load_extern_vdb<float>(device, img, type);
+			else
+#endif
+				file_load_image<TypeDesc::FLOAT, float, float>(device, img, type, texture_limit);
 			break;
 		case IMAGE_DATA_TYPE_BYTE4:
 			file_load_image<TypeDesc::UINT8, uchar, uchar4>(device, img, type, texture_limit);
@@ -1014,8 +1122,10 @@ void ImageManager::device_free_image(Device *, ImageDataType type, int slot)
 
 		if(img->mem) {
 			thread_scoped_lock device_lock(device_mutex);
-			if(img->mem->grid_info){
-				delete img->mem->grid_info;
+			if(img->mem->grid_info && img->mem->grid_type == IMAGE_GRID_TYPE_SPARSE) {
+				device_memory *info = (device_memory*)img->mem->grid_info;
+				delete info;
+				img->mem->grid_info = NULL;
 			}
 			delete img->mem;
 		}
diff --git a/intern/cycles/render/image.h b/intern/cycles/render/image.h
index 75bd9c19e30..e36876da434 100644
--- a/intern/cycles/render/image.h
+++ b/intern/cycles/render/image.h
@@ -145,6 +145,11 @@ private:
 	vector<Image*> images[IMAGE_DATA_NUM_TYPES];
 	void *osl_texture_system;
 
+	bool allocate_sparse_index(Device *device,
+	                           device_memory *tex_img,
+	                           vector<int> *sparse_index,
+	                           string mem_name);
+
 	bool file_load_image_generic(Image *img,
 	                             ImageInput **in,
 	                             int &width,
@@ -152,12 +157,18 @@ private:
 	                             int &depth,
 	                             int &components);
 
-	template<typename StorageType,
-	         typename DeviceType>
+	template<typename DeviceType>
 	void file_load_failed(Image *img,
 	                      ImageDataType type,
 	                      device_vector<DeviceType> *tex_img);
 
+#ifdef WITH_OPENVDB
+	template<typename DeviceType>
+	void file_load_extern_vdb(Device *device,
+	                          Image *img,
+	                          ImageDataType type);
+#endif
+
 	template<TypeDesc::BASETYPE FileFormat,
 	         typename StorageType,
 	         typename DeviceType>
diff --git a/intern/cycles/render/mesh_volume.cpp b/intern/cycles/render/mesh_volume.cpp
index f3138421fd0..49bbd5240c9 100644
--- a/intern/cycles/render/mesh_volume.cpp
+++ b/intern/cycles/render/mesh_volume.cpp
@@ -314,8 +314,9 @@ void VolumeMeshBuilder::convert_quads_to_tris(const vector<QuadData> &quads,
 
 struct VoxelAttributeGrid {
 	void *data;
-	int *grid_info;
+	int *sparse_index;
 	int channels;
+	ImageGridType grid_type;
 };
 
 void MeshManager::create_volume_mesh(Scene *scene,
@@ -330,7 +331,6 @@ void MeshManager::create_volume_mesh(Scene *scene,
 	/* Compute volume parameters. */
 	VolumeParams volume_params;
 	volume_params.resolution = make_int3(0, 0, 0);
-	bool is_openvdb = false;
 	int grid_mem = -1;
 
 	foreach(Attribute& attr, mesh->attributes.attributes) {
@@ -340,7 +340,6 @@ void MeshManager::create_volume_mesh(Scene *scene,
 
 		VoxelAttribute *voxel = attr.data_voxel();
 		device_memory *image_memory = scene->image_manager->image_memory(voxel->slot);
-		device_memory *grid_info = image_memory->grid_info;
 		int3 resolution = make_int3(image_memory->real_width,
 		                            image_memory->real_height,
 		                            image_memory->real_depth);
@@ -355,12 +354,20 @@ void MeshManager::create_volume_mesh(Scene *scene,
 			return;
 		}
 
-		is_openvdb = (image_memory->grid_type == IMAGE_GRID_TYPE_OPENVDB);
-
 		VoxelAttributeGrid voxel_grid;
+
 		voxel_grid.data = image_memory->host_pointer;
 		voxel_grid.channels = image_memory->data_elements;
-		voxel_grid.grid_info = grid_info ? static_cast<int*>(grid_info->host_pointer) : NULL;
+		voxel_grid.grid_type = image_memory->grid_type;
+
+		if(image_memory->grid_type == IMAGE_GRID_TYPE_SPARSE) {
+			device_memory *sparse_mem = (device_memory*)image_memory->grid_info;
+			voxel_grid.sparse_index = static_cast<int*>(sparse_mem->host_pointer);
+		}
+		else {
+			voxel_grid.sparse_index = NULL;
+		}
+
 		voxel_grids.push_back(voxel_grid);
 	}
 
@@ -400,10 +407,6 @@ void MeshManager::create_volume_mesh(Scene *scene,
 	float3 cell_size = make_float3(1.0f/resolution.x,
 	                               1.0f/resolution.y,
 	                               1.0f/resolution.z);
-	const int3 tiled_res = make_int3(get_tile_res(resolution.x),
-	                                 get_tile_res(resolution.y),
-	                                 get_tile_res(resolution.z));
-	const bool using_cuda = (scene->device->info.type == DEVICE_CUDA);
 
 	if(attr) {
 		const Transform *tfm = attr->data_transform();
@@ -424,39 +427,39 @@ void MeshManager::create_volume_mesh(Scene *scene,
 		const VoxelAttributeGrid &voxel_grid = voxel_grids[i];
 		const int channels = voxel_grid.channels;
 
-		if(is_openvdb) {
+		if(voxel_grid.grid_type == IMAGE_GRID_TYPE_OPENVDB) {
 #ifdef WITH_OPENVDB
-			openvdb_build_mesh(&builder, voxel_grid.data, resolution, isovalue, channels > 1);
-#else
-			assert(0);
+			openvdb_build_mesh(&builder, voxel_grid.data, isovalue, channels > 1);
 #endif
 		}
-		else {
-			const int *grid_info = voxel_grid.grid_info;
+		else if(voxel_grid.grid_type == IMAGE_GRID_TYPE_SPARSE) {
 			float *data = static_cast<float*>(voxel_grid.data);
-
 			for(int z = 0; z < resolution.z; ++z) {
 				for(int y = 0; y < resolution.y; ++y) {
 					for(int x = 0; x < resolution.x; ++x) {
-						int voxel_index;
-
-						if(grid_info) {
-							voxel_index = using_cuda ?
-							    compute_index_cuda(grid_info, x, y, z,
-							                       resolution.x, resolution.y, resolution.z,
-							                       tiled_res.x, tiled_res.y, tiled_res.z) :
-							    compute_index(grid_info, x, y, z,
-							                  resolution.x, resolution.y, resolution.z);
-
-							if(voxel_index < 0) {
-								continue;
+						int voxel_index = compute_index(voxel_grid.sparse_index,
+						                                x, y, z,
+						                                resolution.x,
+						                                resolution.y,
+						                                resolution.z) * channels;
+						if(voxel_index >= 0) {
+							for(int c = 0; c < channels; c++) {
+								if(data[voxel_index + c] >= isovalue) {
+									builder.add_node_with_padding(x, y, z);
+									break;
+								}
 							}
 						}
-						else {
-							voxel_index = compute_index(x, y, z, resolution);
-						}
-
-						voxel_index *= channels;
+					}
+				}
+			}
+		}
+		else {
+			float *data = static_cast<float*>(voxel_grid.data);
+			for(int z = 0; z < resolution.z; ++z) {
+				for(int y = 0; y < resolution.y; ++y) {
+					for(int x = 0; x < resolution.x; ++x) {
+						int voxel_index = compute_index(x, y, z, resolution) * channels;
 						for(int c = 0; c < channels; c++) {
 							if(data[voxel_index + c] >= isovalue) {
 								builder.add_node_with_padding(x, y, z);
diff --git a/intern/cycles/render/openvdb.cpp b/intern/cycles/render/openvdb.cpp
index 11bea1242c2..68f2cfa4af0 100644
--- a/intern/cycles/render/openvdb.cpp
+++ b/intern/cycles/render/openvdb.cpp
@@ -17,40 +17,32 @@ struct OpenVDBReader;
 
 CCL_NAMESPACE_BEGIN
 
-/* Comparison and assignment utilities. */
+/* Misc internal helper functions. */
 
-static const bool gte(openvdb::Vec3SGrid::ConstAccessor accessor,
-                      openvdb::math::Coord c, const float f)
+static bool operator >=(const openvdb::math::Vec3s &a, const float &b)
 {
-	openvdb::math::Vec3s v = accessor.getValue(c);
-	return (v.x() >= f || v.y() >= f || v.z() >= f);
+	return a.x() >= b || a.y() >= b || a.z() >= b;
 }
 
-static const bool gte(openvdb::FloatGrid::ConstAccessor accessor,
-                      openvdb::math::Coord c, const float f)
+static void copy(float *des, const openvdb::math::Vec3s *src)
 {
-	return accessor.getValue(c) >= f;
+	*(des + 0) = src->x();
+	*(des + 1) = src->y();
+	*(des + 2) = src->z();
+	*(des + 3) = 1.0f;
 }
 
-static void copy(openvdb::Vec3SGrid::ConstAccessor accessor,
-                 openvdb::math::Coord c, float *f)
+static void copy(float *des, const float *src)
 {
-	openvdb::math::Vec3s v = accessor.getValue(c);
-	*(f + 0) = v.x();
-	*(f + 1) = v.y();
-	*(f + 2) = v.z();
-	*(f + 3) = 1.0f;
+	*des = *src;
 }
 
-static void copy(openvdb::FloatGrid::ConstAccessor accessor,
-                 openvdb::math::Coord c, float *f)
+static const int tile_index(openvdb::math::Coord start, int3 tiled_res)
 {
-	*f = accessor.getValue(c);
+	return compute_index(start.x() / TILE_SIZE, start.y() / TILE_SIZE,
+	                     start.z() / TILE_SIZE, tiled_res.x, tiled_res.y);
 }
 
-/* Misc internal helper functions.
- * Logging should be done by callers. */
-
 /* Simple range shift for grids with non-zero background values. May have
  * strange results depending on the grid. */
 static void shift_range(openvdb::Vec3SGrid::Ptr grid)
@@ -86,8 +78,8 @@ template<typename GridType>
 static bool get_grid(const string& filepath,
                      const string& grid_name,
                      typename GridType::Ptr& grid,
-                     typename GridType::ConstAccessor& accessor,
-                     int3& resolution)
+                     int3 *resolution,
+                     openvdb::math::Coord *minimum_bound)
 {
 	using namespace openvdb;
 
@@ -103,33 +95,35 @@ static bool get_grid(const string& filepath,
 		return false;
 	}
 
-	int min_bound[3], res[3];
-	OpenVDBReader_get_bounds(reader, min_bound, NULL, res, NULL, NULL, NULL);
-
-	/* In order to keep sampling uniform, we expect a volume's bound to begin at
-	 * (0, 0, 0) in object space. External VDBs may have a non-zero origin, so
-	 * all voxels must be translated. This process may be memory inefficient. */
+	grid = gridPtrCast<GridType>(reader->getGrid(grid_name));
 
-	typename GridType::Ptr orig_grid = gridPtrCast<GridType>(reader->getGrid(grid_name));
+	/* Verify that leaf dimensions match internal tile dimensions. */
+	typename GridType::TreeType::LeafCIter iter = grid->tree().cbeginLeaf();
+	if(iter) {
+		const math::Coord dim = iter.getLeaf()->getNodeBoundingBox().dim();
 
-	if(min_bound[0] == 0 && min_bound[1] == 0 && min_bound[2] == 0) {
-		grid = orig_grid;
-	}
-	else {
-		grid = GridType::create();
-		math::Mat4d xform = math::Mat4d::identity();
-		math::Vec3d translation(-min_bound[0], -min_bound[1], -min_bound[2]);
-		xform.setTranslation(translation);
-		tools::GridTransformer transformer(xform);
-		transformer.transformGrid<tools::PointSampler, GridType>(*orig_grid, *grid);
+		if(dim[0] != TILE_SIZE || dim[1] != TILE_SIZE || dim[2] != TILE_SIZE) {
+			VLOG(1) << "Cannot load grid " << grid->getName() << " from "
+			        << filepath << ", leaf dimensions are "
+			        << dim[0] << "x" << dim[1] << "x" << dim[2];
+			OpenVDBReader_free(reader);
+			return false;
+		}
 	}
 
 	/* Need to account for external grids with a non-zero background value and
 	 * voxels below background value. */
 	shift_range(grid);
 
-	accessor = grid->getConstAccessor();
-	resolution = make_int3(res[0], res[1], res[2]);
+	int min_bound[3], res[3];
+	OpenVDBReader_get_bounds(reader, min_bound, NULL, res, NULL, NULL, NULL);
+
+	if(resolution) {
+		*resolution = make_int3(res[0], res[1], res[2]);
+	}
+	if(minimum_bound) {
+		*minimum_bound = math::Coord(min_bound[0], min_bound[1], min_bound[2]);
+	}
 
 	OpenVDBReader_free(reader);
 	return true;
@@ -182,17 +176,34 @@ static device_memory *openvdb_load_device_extern(Device *device,
 	using namespace openvdb;
 
 	typename GridType::Ptr grid = GridType::create();
-	typename GridType::ConstAccessor accessor = grid->getConstAccessor();
+	typename GridType::Ptr orig_grid = GridType::create();
 	int3 resolution;
+	math::Coord min_bound;
 
-	if(!get_grid<GridType>(filepath, grid_name, grid, accessor, resolution)) {
+	if(!get_grid<GridType>(filepath, grid_name, orig_grid, &resolution, &min_bound)) {
 		return NULL;
 	}
 
+	/* In order to keep sampling uniform, we expect a volume's bound to begin at
+	 * (0, 0, 0) in object space. External VDBs may have a non-zero origin, so
+	 * all voxels must be translated. This process may be memory inefficient. */
+
+	if(min_bound.x() == 0 && min_bound.y() == 0 && min_bound.z() == 0) {
+		grid = orig_grid;
+	}
+	else {
+		math::Mat4d xform = math::Mat4d::identity();
+		math::Vec3d translation(-min_bound.x(), -min_bound.y(), -min_bound.z());
+		xform.setTranslation(translation);
+		tools::GridTransformer transformer(xform);
+		transformer.transformGrid<tools::PointSampler, GridType>(*orig_grid, *grid);
+	}
+
+	typename GridType::ConstAccessor accessor = grid->getConstAccessor();
+
 	device_openvdb<GridType> *tex_img =
-			new device_openvdb<GridType>(device, mem_name.c_str(),
-	                                     MEM_TEXTURE, grid, accessor,
-	                                     resolution);
+			new device_openvdb<GridType>(device, mem_name.c_str(), MEM_TEXTURE,
+	                                     grid, accessor, resolution);
 
 	tex_img->interpolation = interpolation;
 	tex_img->extension = extension;
@@ -222,189 +233,203 @@ device_memory *openvdb_load_device_extern(Device *device,
 	}
 }
 
-/* Load OpenVDB file to sparse grid. Based on util/util_sparse_grid.h */
-template<typename GridType>
-static const bool openvdb_check_tile_active(typename GridType::ConstAccessor accessor,
-                                            int x, int y, int z,
-                                            float threshold, int3 resolution)
-{
-	using namespace openvdb;
-
-	math::Coord ijk;
-	int &i = ijk[0], &j = ijk[1], &k = ijk[2];
-
-	const int max_i = min(x + TILE_SIZE, resolution.x);
-	const int max_j = min(y + TILE_SIZE, resolution.y);
-	const int max_k = min(z + TILE_SIZE, resolution.z);
-
-	for(k = z; k < max_k; ++k) {
-		for(j = y; j < max_j; ++j) {
-			for(i = x; i < max_i; ++i) {
-				if(gte(accessor, ijk, threshold)) {
-					return true;
-				}
-			}
-		}
-	}
-	return false;
-}
-
-template<typename GridType>
-static bool openvdb_load_sparse(const string& filepath,
-                                const string& grid_name,
-                                const int channels,
-                                const float threshold,
-                                vector<float> *sparse_grid,
-		                        vector<int> *grid_info)
+/* Load OpenVDB file to texture grid. */
+template<typename GridType, typename T>
+static void openvdb_load_preprocess(const string& filepath,
+                                    const string& grid_name,
+                                    const int channels,
+                                    const float threshold,
+                                    vector<int> *sparse_index,
+                                    int &sparse_size)
 {
 	using namespace openvdb;
 
 	typename GridType::Ptr grid = GridType::create();
-	typename GridType::ConstAccessor accessor = grid->getConstAccessor();
 	int3 resolution;
+	math::Coord min_bound;
 
-	if(!get_grid<GridType>(filepath, grid_name, grid, accessor, resolution)) {
-		return false;
+	if(!get_grid<GridType>(filepath, grid_name, grid, &resolution, &min_bound) ||
+	   !(channels == 4 || channels == 1))
+	{
+		return;
 	}
 
-	const int tile_count = get_tile_res(resolution.x) *
-	                       get_tile_res(resolution.y) *
-	                       get_tile_res(resolution.z);
+	const int3 tiled_res = make_int3(get_tile_res(resolution.x),
+	                                 get_tile_res(resolution.y),
+	                                 get_tile_res(resolution.z));
+	const int3 last_tile = make_int3(resolution.x % TILE_SIZE,
+	                                 resolution.y % TILE_SIZE,
+	                                 resolution.z % TILE_SIZE);
+	const int tile_count = tiled_res.x * tiled_res.y * tiled_res.z;
 	const int tile_pix_count = TILE_SIZE * TILE_SIZE * TILE_SIZE * channels;
 
-	/* Initial prepass to find active tiles. */
-	grid_info->resize(tile_count);
-	int tile = 0, active_count = 0;
-
-	for(int z = 0; z < resolution.z; z += TILE_SIZE) {
-		for(int y = 0; y < resolution.y; y += TILE_SIZE) {
-			for(int x = 0; x < resolution.x; x += TILE_SIZE, ++tile) {
-				int is_active = openvdb_check_tile_active<GridType>(accessor,
-				                                                    x, y, z,
-				                                                    threshold,
-				                                                    resolution);
-				active_count += is_active;
-				/* 0 if active, -1 if inactive. */
-				grid_info->at(tile) = is_active - 1;
+	sparse_index->resize(tile_count, -1); /* 0 if active, -1 if inactive. */
+	int voxel_count = 0;
+
+	for (typename GridType::TreeType::LeafCIter iter = grid->tree().cbeginLeaf(); iter; ++iter) {
+		const typename GridType::TreeType::LeafNodeType *leaf = iter.getLeaf();
+		const T *data = leaf->buffer().data();
+		const math::Coord start = leaf->getNodeBoundingBox().getStart() - min_bound;
+
+		for(int i = 0; i < tile_pix_count; ++i) {
+			if(data[i] >= threshold) {
+				sparse_index->at(tile_index(start, tiled_res)) = 0;
+
+				/* Calculate how many voxels are in this tile. */
+				int tile_width = (start.x() + TILE_SIZE > resolution.x) ? last_tile.x : TILE_SIZE;
+				int tile_height = (start.y() + TILE_SIZE > resolution.y) ? last_tile.y : TILE_SIZE;
+				int tile_depth = (start.z() + TILE_SIZE > resolution.z) ? last_tile.z : TILE_SIZE;
+				voxel_count += tile_width * tile_height * tile_depth;
+
+				break;
 			}
 		}
 	}
 
 	/* Check memory savings. */
-	int sparse_mem_use = (tile_count * sizeof(int) +
-	                      active_count * tile_pix_count * sizeof(float));
-	int dense_mem_use = resolution.x * resolution.y *
-	                    resolution.z * channels * sizeof(float);
-
-	if(sparse_mem_use >= dense_mem_use) {
+	const int sparse_mem_use = tile_count * sizeof(int) + voxel_count * channels * sizeof(float);
+	const int dense_mem_use = resolution.x * resolution.y * resolution.z * channels * sizeof(float);
+
+	if(sparse_mem_use < dense_mem_use) {
+		VLOG(1) << "Memory of " << grid_name << " decreased from "
+				<< string_human_readable_size(dense_mem_use) << " to "
+				<< string_human_readable_size(sparse_mem_use);
+		sparse_size = voxel_count * channels;
+	}
+	else {
 		VLOG(1) << "Memory of " << grid_name << " increased from "
 		        << string_human_readable_size(dense_mem_use) << " to "
 		        << string_human_readable_size(sparse_mem_use)
 		        << ", not using sparse grid";
-		return false;
+		sparse_size = -1;
+		sparse_index->resize(0);
 	}
-
-	VLOG(1) << "Memory of " << grid_name << " decreased from "
-			<< string_human_readable_size(dense_mem_use) << " to "
-			<< string_human_readable_size(sparse_mem_use);
-
-	/* Populate the sparse grid. */
-	sparse_grid->resize(active_count * tile_pix_count);
-	float *sg = &(*sparse_grid)[0];
-
-	int voxel = 0;
-	tile = 0;
-
-	for(int z = 0; z < resolution.z; z += TILE_SIZE) {
-		for(int y = 0; y < resolution.y; y += TILE_SIZE) {
-			for(int x = 0; x < resolution.x; x += TILE_SIZE, ++tile) {
-				if(grid_info->at(tile) == -1) {
-					continue;
-				}
-
-				grid_info->at(tile) = voxel / channels;
-
-				/* Populate the tile. */
-				const int max_i = min(x + TILE_SIZE, resolution.x);
-				const int max_j = min(y + TILE_SIZE, resolution.y);
-				const int max_k = min(z + TILE_SIZE, resolution.z);
-
-				math::Coord ijk;
-				int &i = ijk[0], &j = ijk[1], &k = ijk[2];
-
-				for(k = z; k < max_k; ++k) {
-					for(j = y; j < max_j; ++j) {
-						for(i = x; i < max_i; ++i, ++voxel) {
-							copy(accessor, ijk, sg + voxel);
-						}
-					}
-				}
-			}
-		}
-	}
-
-	return true;
 }
 
-bool openvdb_load_sparse(const string& filepath,
-                         const string& grid_name,
-                         const int channels,
-                         const float threshold,
-                         vector<float> *sparse_grid,
-                         vector<int> *grid_info)
+void openvdb_load_preprocess(const string& filepath,
+                             const string& grid_name,
+                             const int channels,
+                             const float threshold,
+                             vector<int> *sparse_index,
+                             int &sparse_size)
 {
 	if(channels > 1) {
-		return openvdb_load_sparse<openvdb::Vec3SGrid>(filepath, grid_name,
-		                                               channels, threshold,
-		                                               sparse_grid, grid_info);
+		return openvdb_load_preprocess<openvdb::Vec3SGrid, openvdb::math::Vec3s>(
+		            filepath, grid_name, channels, threshold, sparse_index, sparse_size);
 	}
 	else {
-		return openvdb_load_sparse<openvdb::FloatGrid>(filepath, grid_name,
-		                                               channels, threshold,
-		                                               sparse_grid, grid_info);
+		return openvdb_load_preprocess<openvdb::FloatGrid, float>(
+		            filepath, grid_name, channels, threshold, sparse_index, sparse_size);
 	}
 }
 
-/* Load OpenVDB file to dense grid. */
-template<typename GridType>
-static bool openvdb_load_dense(const string& filepath, const string& grid_name,
-	                           float *data, const int channels)
+
+template<typename GridType, typename T>
+static void openvdb_load_image(const string& filepath,
+                               const string& grid_name,
+                               const int channels,
+                               float *image,
+                               vector<int> *sparse_index)
 {
 	using namespace openvdb;
 
 	typename GridType::Ptr grid = GridType::create();
-	typename GridType::ConstAccessor accessor = grid->getConstAccessor();
 	int3 resolution;
+	math::Coord min_bound;
 
-	if(!get_grid<GridType>(filepath, grid_name, grid, accessor, resolution)) {
-		return false;
+	if(!get_grid<GridType>(filepath, grid_name, grid, &resolution, &min_bound) ||
+	   !(channels == 4 || channels == 1))
+	{
+		return;
 	}
 
-	math::Coord xyz;
-	int &x = xyz[0], &y = xyz[1], &z = xyz[2];
-	int index = 0;
+	bool make_sparse = false;
+	if(sparse_index) {
+		if(sparse_index->size() > 0) {
+			make_sparse = true;
+		}
+	}
 
-	for (z = 0; z < resolution.z; ++z) {
-		for (y = 0; y < resolution.y; ++y) {
-			for (x = 0; x < resolution.x; ++x, index += channels) {
-				copy(accessor, xyz, data + index);
+	if(make_sparse) {
+		/* Load VDB as sparse image. */
+		const int3 tiled_res = make_int3(get_tile_res(resolution.x),
+		                                 get_tile_res(resolution.y),
+		                                 get_tile_res(resolution.z));
+		const int3 last_tile = make_int3(resolution.x % TILE_SIZE,
+		                                 resolution.y % TILE_SIZE,
+		                                 resolution.z % TILE_SIZE);
+		int start_index = 0;
+
+		for (typename GridType::TreeType::LeafCIter iter = grid->tree().cbeginLeaf(); iter; ++iter) {
+			const typename GridType::TreeType::LeafNodeType *leaf = iter.getLeaf();
+
+			const math::Coord start = leaf->getNodeBoundingBox().getStart() - min_bound;
+			int tile = tile_index(start, tiled_res);
+			if(sparse_index->at(tile) == -1) {
+				continue;
+			}
+			sparse_index->at(tile) = start_index / channels;
+
+
+			const T *vdb_tile = leaf->buffer().data();
+			float *arr_tile = image + start_index;
+
+
+			const int tile_width = (start.x() + TILE_SIZE > resolution.x) ? last_tile.x : TILE_SIZE;
+			const int tile_height = (start.y() + TILE_SIZE > resolution.y) ? last_tile.y : TILE_SIZE;
+			const int tile_depth = (start.z() + TILE_SIZE > resolution.z) ? last_tile.z : TILE_SIZE;
+
+			/* Index computation by coordinates is reversed in VDB grids. */
+			for(int k = 0; k < tile_depth; ++k) {
+				for(int j = 0; j < tile_height; ++j) {
+					for(int i = 0; i < tile_width; ++i) {
+						int arr_index = compute_index(i, j, k, tile_width, tile_height);
+						int vdb_index = compute_index(k, j, i, TILE_SIZE, TILE_SIZE);
+						copy(arr_tile + arr_index, vdb_tile + vdb_index);
+					}
+				}
+			}
+
+			start_index += tile_width * tile_height * tile_depth;
+		}
+	}
+	else {
+		/* Load VDB as dense image. */
+		for (typename GridType::TreeType::LeafCIter iter = grid->tree().cbeginLeaf(); iter; ++iter) {
+			const typename GridType::TreeType::LeafNodeType *leaf = iter.getLeaf();
+			const T *vdb_tile = leaf->buffer().data();
+			const math::Coord start = leaf->getNodeBoundingBox().getStart() - min_bound;
+
+			for (int k = 0; k < TILE_SIZE; ++k) {
+				for (int j = 0; j < TILE_SIZE; ++j) {
+					for (int i = 0; i < TILE_SIZE; ++i) {
+						int arr_index = compute_index(start.x() + i,
+													  start.y() + j,
+													  start.z() + k,
+													  resolution.x,
+													  resolution.y);
+						int vdb_index = compute_index(k, j, i, TILE_SIZE, TILE_SIZE);
+						copy(image + arr_index, vdb_tile + vdb_index);
+					}
+				}
 			}
 		}
 	}
-
-	return true;
 }
 
-bool openvdb_load_dense(const string& filepath, const string& grid_name,
-                        float *data, const int channels)
+void openvdb_load_image(const string& filepath,
+                        const string& grid_name,
+                        const int channels,
+                        float *image,
+                        vector<int> *sparse_index)
 {
 	if(channels > 1) {
-		return openvdb_load_dense<openvdb::Vec3SGrid>(filepath, grid_name,
-		                                              data, channels);
+		return openvdb_load_image<openvdb::Vec3SGrid, openvdb::math::Vec3s>(
+		    filepath, grid_name, channels, image, sparse_index);
 	}
 	else {
-		return openvdb_load_dense<openvdb::FloatGrid>(filepath, grid_name,
-		                                              data, channels);
+		return openvdb_load_image<openvdb::FloatGrid, float>(
+		    filepath, grid_name, channels, image, sparse_index);
 	}
 }
 
@@ -510,39 +535,44 @@ device_memory *openvdb_load_device_intern(Device *device,
 /* Volume Mesh Builder functions. */
 
 template<typename GridType>
-static void openvdb_build_mesh(VolumeMeshBuilder *builder, void *v_accessor,
-                               const int3 resolution, const float threshold)
+static void openvdb_build_mesh(VolumeMeshBuilder *builder,
+                               void *v_grid,
+                               const float threshold)
 {
 	using namespace openvdb;
 
-	typename GridType::ConstAccessor *acc =
-	        static_cast<typename GridType::ConstAccessor*>(v_accessor);
-
-	math::Coord xyz;
-	int &x = xyz[0], &y = xyz[1], &z = xyz[2];
-
-	for (z = 0; z < resolution.z; ++z) {
-		for (y = 0; y < resolution.y; ++y) {
-			for (x = 0; x < resolution.x; ++x) {
-				if(gte(*acc, xyz, threshold)) {
-					builder->add_node_with_padding(x, y, z);
+	typename GridType::Ptr grid = *static_cast<typename GridType::Ptr*>(v_grid);
+
+	for (typename GridType::TreeType::LeafCIter iter = grid->tree().cbeginLeaf(); iter; ++iter) {
+		const typename GridType::TreeType::LeafNodeType *leaf = iter.getLeaf();
+		const float *data = (float*)leaf->buffer().data();
+		const math::Coord start = leaf->getNodeBoundingBox().getStart();
+
+		for (int k = 0; k < TILE_SIZE; ++k) {
+			for (int j = 0; j < TILE_SIZE; ++j) {
+				for (int i = 0; i < TILE_SIZE; ++i) {
+					int vdb_index = compute_index(k, j, i, TILE_SIZE, TILE_SIZE);
+					if(data[vdb_index] >= threshold) {
+						builder->add_node_with_padding(start.x() + i,
+						                               start.y() + j,
+						                               start.z() + k);
+					}
 				}
 			}
 		}
 	}
 }
 
-void openvdb_build_mesh(VolumeMeshBuilder *builder, void *v_accessor,
-                        const int3 resolution, const float threshold,
+void openvdb_build_mesh(VolumeMeshBuilder *builder,
+                        void *v_grid,
+                        const float threshold,
                         const bool is_vec)
 {
 	if(is_vec) {
-		openvdb_build_mesh<openvdb::Vec3SGrid>(builder, v_accessor,
-		                                       resolution, threshold);
+		openvdb_build_mesh<openvdb::Vec3SGrid>(builder, v_grid, threshold);
 	}
 	else {
-		openvdb_build_mesh<openvdb::FloatGrid>(builder, v_accessor,
-		                                       resolution, threshold);
+		openvdb_build_mesh<openvdb::FloatGrid>(builder, v_grid, threshold);
 	}
 }
 
diff --git a/intern/cycles/render/openvdb.h b/intern/cycles/render/openvdb.h
index 06e404059d6..37657800372 100644
--- a/intern/cycles/render/openvdb.h
+++ b/intern/cycles/render/openvdb.h
@@ -20,15 +20,18 @@ device_memory *openvdb_load_device_extern(Device *device,
                                           const ExtensionType& extension,
                                           const bool is_vec);
 
-bool openvdb_load_sparse(const string& filepath,
-                         const string& grid_name,
-                         const int channels,
-                         const float threshold,
-                         vector<float> *sparse_grid,
-                         vector<int> *grid_info);
-
-bool openvdb_load_dense(const string& filepath, const string& grid_name,
-                        float *data, const int channels);
+void openvdb_load_preprocess(const string& filepath,
+                             const string& grid_name,
+                             const int channels,
+                             const float threshold,
+                             vector<int> *sparse_index,
+                             int &sparse_size);
+
+void openvdb_load_image(const string& filepath,
+                        const string& grid_name,
+                        const int channels,
+                        float *image,
+                        vector<int> *sparse_index);
 
 device_memory *openvdb_load_device_intern(Device *device,
                                           const float *data,
@@ -38,8 +41,9 @@ device_memory *openvdb_load_device_intern(Device *device,
                                           const ExtensionType& extension,
                                           const bool is_vec);
 
-void openvdb_build_mesh(VolumeMeshBuilder *builder, void *v_accessor,
-                        const int3 resolution, const float threshold,
+void openvdb_build_mesh(VolumeMeshBuilder *builder,
+                        void *v_grid,
+                        const float threshold,
                         const bool is_vec);
 
 CCL_NAMESPACE_END
diff --git a/intern/cycles/util/util_sparse_grid.h b/intern/cycles/util/util_sparse_grid.h
index 8bc90e8b504..e179c905f6b 100644
--- a/intern/cycles/util/util_sparse_grid.h
+++ b/intern/cycles/util/util_sparse_grid.h
@@ -137,34 +137,6 @@ const inline int compute_index_cuda(const int *grid_info,
 }
 
 template<typename T>
-static const bool check_tile_active(const T *dense_grid,
-                                    int x, int y, int z,
-                                    const T threshold,
-                                    const int width,
-                                    const int height,
-                                    const int depth,
-                                    const int channels)
-{
-	const int max_i = min(x + TILE_SIZE, width);
-	const int max_j = min(y + TILE_SIZE, height);
-	const int max_k = min(z + TILE_SIZE, depth);
-
-	for(int k = z; k < max_k; ++k) {
-		for(int j = y; j < max_j; ++j) {
-			for(int i = x; i < max_i; ++i) {
-				int index = compute_index(i, j, k, width, height) * channels;
-				for(int c = 0; c < channels; ++c) {
-					if(dense_grid[index + c] >= threshold) {
-						return true;
-					}
-				}
-			}
-		}
-	}
-	return false;
-}
-
-template<typename T>
 bool create_sparse_grid(const T *dense_grid,
                         const int width,
                         const int height,
@@ -181,27 +153,44 @@ bool create_sparse_grid(const T *dense_grid,
 	const int tile_count = get_tile_res(width) *
 	                       get_tile_res(height) *
 	                       get_tile_res(depth);
-	const int tile_pix_count = TILE_SIZE * TILE_SIZE * TILE_SIZE * channels;
 
 	/* Initial prepass to find active tiles. */
-	grid_info->resize(tile_count);
-	int tile = 0, active_count = 0;
+	grid_info->resize(tile_count, -1);  /* 0 if active, -1 if inactive. */
+	int tile = 0, voxel = 0, voxel_count = 0;
 
 	for(int z = 0; z < depth; z += TILE_SIZE) {
 		for(int y = 0; y < height; y += TILE_SIZE) {
 			for(int x = 0; x < width; x += TILE_SIZE, ++tile) {
-				int is_active = check_tile_active<T>(dense_grid, x, y, z, threshold,
-				                                     width, height, depth, channels);
-				active_count += is_active;
-				/* 0 if active, -1 if inactive. */
-				grid_info->at(tile) = is_active - 1;
+				bool is_active = false;
+				const int max_i = min(x + TILE_SIZE, width);
+				const int max_j = min(y + TILE_SIZE, height);
+				const int max_k = min(z + TILE_SIZE, depth);
+
+				voxel = 0;
+				for(int k = z; k < max_k; ++k) {
+					for(int j = y; j < max_j; ++j) {
+						for(int i = x; i < max_i; ++i, ++voxel) {
+							int index = compute_index(i, j, k, width, height) * channels;
+							for(int c = 0; c < channels; ++c) {
+								if(dense_grid[index + c] >= threshold) {
+									is_active = true;
+									break;
+								}
+							}
+						}
+					}
+				}
+
+				if(is_active) {
+					grid_info->at(tile) = 0;
+					voxel_count += voxel;
+				}
 			}
 		}
 	}
 
 	/* Check memory savings. */
-	int sparse_mem_use = (tile_count * sizeof(int) +
-	                      active_count * tile_pix_count * sizeof(T));
+	int sparse_mem_use = tile_count * sizeof(int) + voxel_count * channels * sizeof(T);
 	int dense_mem_use = width * height * depth * channels * sizeof(T);
 
 	if(sparse_mem_use >= dense_mem_use) {
@@ -217,10 +206,8 @@ bool create_sparse_grid(const T *dense_grid,
 			<< string_human_readable_size(sparse_mem_use);
 
 	/* Populate the sparse grid. */
-	sparse_grid->resize(active_count * tile_pix_count);
-
-	int voxel = 0;
-	tile = 0;
+	sparse_grid->resize(voxel_count * channels);
+	voxel = tile = 0;
 
 	for(int z = 0; z < depth; z += TILE_SIZE) {
 		for(int y = 0; y < height; y += TILE_SIZE) {