 static __device__ void Reside (struct particle *, unsigned int *);
 static __device__ void Swap (struct particle *, struct particle *);
 
+static __device__ int blocksComplete = 0;
+
 extern "C" __global__ void ParticleSort (unsigned int *global_mem,
 					 unsigned int *transblock_buffers,
 					 unsigned int *buffer_flags,

	}
	switch (role) {
	case BEGINNING:
		*here = 0;
		// fall through
	case MIDDLE:


extern "C" void sort (unsigned int *buffer, unsigned long size)
{
	dim3 grid ((size - 1) / BLOCK + 1);
	dim3 block (BLOCK);
	size_t global_mem_size = size * sizeof(int);
	size_t transblock_buffers_size = BUFFER_SIZE * (block.x - 1) * sizeof(int);
