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

				atomicAdd(&blocksComplete, -1);
			}
		}
		if (++i & 0xF == 0) {
			switch (buff_role) {
			case LBUFF:
				while (buffer_flags[(blockIdx.x-1)*2+1]); /*busy wait*/ 
				left_incoming[threadIdx.x] = transblock_buffers[(blockIdx.x-1)*_2XBUFFER_SIZE+threadIdx.x];
				transblock_buffers[(blockIdx.x-1)*_2XBUFFER_SIZE+threadIdx.x] = 0;
				transblock_buffers[(blockIdx.x - 1) * _2XBUFFER_SIZE + BUFFER_SIZE + threadIdx.x] = left_outgoing[threadIdx.x];
				atomicExch(buffer_flags + (blockIdx.x - 1) * 2, 0);
				atomicExch(buffer_flags + (blockIdx.x - 1) * 2 + 1, 1);
				break;
			case RBUFF:
				while (buffer_flags[blockIdx.x*2]); /*busy wait*/
				transblock_buffers[blockIdx.x*_2XBUFFER_SIZE+threadIdx.x-blockDim.x+BUFFER_SIZE] = right_outgoing[threadIdx.x-blockDim.x+BUFFER_SIZE];
				right_incoming[threadIdx.x-blockDim.x+BUFFER_SIZE] =
					transblock_buffers[blockIdx.x*_2XBUFFER_SIZE+BUFFER_SIZE+threadIdx.x-blockDim.x+BUFFER_SIZE];

	dim3 grid ((size - 1) / BLOCK + 1);
	dim3 block (BLOCK);
	size_t global_mem_size = size * sizeof(int);
	size_t transblock_buffers_size = _2XBUFFER_SIZE * (grid.x - 1) * sizeof(int);
	size_t buffer_flags_size = (grid.x - 1) * 2 * sizeof(int);

	ErrorCheck(cudaMalloc(&global_mem, global_mem_size), "cudaMalloc global");
	ErrorCheck(cudaMemcpy(global_mem, buffer, global_mem_size, cudaMemcpyHostToDevice),
