 //	return res;
 //}
 
-MC_GLOBAL void spParticleCountKernel(dim3 dim, spParticlePage **buckets, int * res)
-{
-
+MC_GLOBAL void spParticleCountKernel(dim3 dim, dim3 offset, dim3 count, spParticlePage **buckets,
+		int * page_count_device)
+{
 	dim3 idx = spParallelBlockIdx();
 
-	if (dim.x > idx.x && dim.y > idx.y && dim.z > idx.z)
-	{
-		size_type pos = idx.x + (idx.y + idx.z * dim.y) * dim.y;
+	if (count.x > idx.x && count.y > idx.y && count.z > idx.z)
+	{
+		size_type pos = idx.x + offset.x + (idx.y + offset.y + (idx.z + offset.z) * dim.y) * dim.x;
+		size_type pos2 = idx.x + (idx.y + idx.z * count.y) * count.x;
 
 		int count = 0;
 


	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;
}

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