 
 extern "C" void sort (unsigned int *buffer, unsigned long size)
 {
-	dim3 grid (1);
+	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);

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

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,
