 
 	spIOStreamPWD(os, curr_path);
 	spIOStreamOpen(os, new_path);
-
-	size_type number_of_cell = spMeshGetNumberOfEntity(sp->m, sp->iform/*volume*/);
-
-	int *page_count_device;
-	spParallelDeviceMalloc((void**) (&page_count_device), number_of_cell * sizeof(int));
-
-	spParticleCountKernel<<<sp->m->dims,1>>>( sp->m->dims, sp->m_buckets_, page_count_device);
-
-	size_type num_of_pages = 0;
-	int *page_count_host;
-
-	spParallelHostMalloc((void**) (&page_count_host), number_of_cell * sizeof(int));
-
-	spParallelMemcpy((void*) (page_count_host), (void*) (page_count_device), number_of_cell * sizeof(int));
-
-	for (int i = 0; i < number_of_cell; ++i)
-	{
-		size_type old = num_of_pages;
-		num_of_pages += page_count_host[i];
-		page_count_host[i] = old;
-	}
-	spParallelMemcpy((void*) (page_count_device), (void*) (page_count_host), number_of_cell * sizeof(int));
-
-	spParallelHostFree((void**) &page_count_host);
-	/*****************************************************************************************/
-
-	void** page_data_ptr_device;
-	void** page_data_ptr_host;
-	MeshEntityId * page_id_device;
-	MeshEntityId * page_id_host;
-
-	spParallelDeviceMalloc((void**) (&page_data_ptr_device), sp->number_of_pages * sizeof(spPage*));
-	spParallelDeviceMalloc((void**) (&page_id_device), sp->number_of_pages * sizeof(MeshEntityId));
-
-	spParticleDumpKernel<<<sp->m->dims,1>>>( page_count_device ,sp->m_buckets_,page_data_ptr_device, page_id_device, page_count_device);
-
-	spParallelDeviceFree((void**) &page_count_device);
-
-	spParallelHostMalloc((void**) (&page_data_ptr_host), sp->number_of_pages * sizeof(spPage*));
-	spParallelMemcpy((void*) (page_data_ptr_host), (void*) (page_data_ptr_device),
-			sp->number_of_pages * sizeof(spPage*));
-	spParallelDeviceFree((void**) &page_data_ptr_device);
-
-	spParallelHostMalloc((void**) (&page_id_host), sp->number_of_pages * sizeof(MeshEntityId));
-	spParallelMemcpy((void*) (page_id_host), (void*) (page_id_device), sp->number_of_pages * sizeof(MeshEntityId));
-	spParallelDeviceFree((void**) &page_id_device);
-
-	size_type dims[2];
-
-	dims[0] = num_of_pages;
-	dims[1] = SP_NUMBER_OF_ENTITIES_IN_PAGE;
-
-	spIOStreamWriteSimple(os, "MeshId", (int) SP_TYPE_MeshEntityId, page_id_host, 1, dims, NULL, dims, (int) flag);
+	dim3 count;
+	count.x = sp->m->i_upper.x - sp->m->i_lower.x;
+	count.y = sp->m->i_upper.y - sp->m->i_lower.y;
+	count.z = sp->m->i_upper.z - sp->m->i_lower.z;
+	num_of_pages = spParticleCopyOut(sp, sp->m->i_lower, count, &page_id_host, &page_data_ptr_host);
+
+	size_type f_dims[2];
+
+	f_dims[0] = num_of_pages;
+	f_dims[1] = SP_NUMBER_OF_ENTITIES_IN_PAGE;
+
+	spIOStreamWriteSimple(os, "MeshId", (int) SP_TYPE_MeshEntityId, page_id_host, 1, f_dims, NULL, f_dims, (int) flag);
 
 	for (int i = 0; i < sp->num_of_attrs; ++i)
 	{


	spParallelMemset((void *) ((sp->m_buckets_)), 0x0, sizeof(spPage*) * number_of_cell);

	spParticleDeployKernel<<<sp->m->dims,NUMBER_OF_THREADS_PER_BLOCK>>>(sp->m_pages_, sp->number_of_pages,sp->page_size_in_byte );

	spParallelDeviceSync();        //wait for iteration to finish


		++pos;
		pg = pg->next;
	}
}
size_type spParticleCopyOut(spParticle const *sp, dim3 offset, dim3 count, MeshEntityId ** page_id_host,
		void*** page_data_ptr_host)
{
	size_type num_of_pages = 0;

	size_type number_of_cell = (count.x) * (count.y) * (count.z);

	int *page_count_device;

	spParallelDeviceMalloc((void**) (&page_count_device), number_of_cell * sizeof(int));

	spParticleCountKernel<<<count, 1>>>(sp->m->dims, offset, count, sp->m_buckets_, page_count_device);

	int *page_count_host;

	spParallelHostMalloc((void**) (&page_count_host), number_of_cell * sizeof(int));

	spParallelMemcpy((void*) (page_count_host), (void*) (page_count_device), number_of_cell * sizeof(int));

	for (int i = 0; i < number_of_cell; ++i)
	{
		size_type old = num_of_pages;
		num_of_pages += page_count_host[i];
		page_count_host[i] = old;
	}
	spParallelMemcpy((void*) (page_count_device), (void*) (page_count_host), number_of_cell * sizeof(int));

	spParallelHostFree((void**) &page_count_host);
	/*****************************************************************************************/

	void** page_data_ptr_device;

	MeshEntityId * page_id_device;

	spParallelDeviceMalloc((void**) (&page_data_ptr_device), sp->number_of_pages * sizeof(spPage*));
	spParallelDeviceMalloc((void**) (&page_id_device), sp->number_of_pages * sizeof(MeshEntityId));

	spParticleDumpKernel<<<sp->m->dims, 1>>>(page_count_device, sp->m_buckets_, page_data_ptr_device, page_id_device);

	spParallelDeviceFree((void**) &page_count_device);

	spParallelHostMalloc((void**) (page_data_ptr_host), sp->number_of_pages * sizeof(spPage*));
	spParallelMemcpy((void*) (*page_data_ptr_host), (void*) (page_data_ptr_device),
			sp->number_of_pages * sizeof(spPage*));
	spParallelDeviceFree((void**) &page_data_ptr_device);

	spParallelHostMalloc((void**) (page_id_host), sp->number_of_pages * sizeof(MeshEntityId));
	spParallelMemcpy((void*) (*page_id_host), (void*) (page_id_device), sp->number_of_pages * sizeof(MeshEntityId));
	spParallelDeviceFree((void**) &page_id_device);

	return num_of_pages;
}

/*****************************************************************************************/
