 				atomicAdd(&blocksComplete, -1);
 			}
 		}
-		if (++i & 0xF == 0) {
+		if ((++i & 0xF) == 0) {
 			switch (buff_role) {
 			case LBUFF:
-				while (buffer_flags[(blockIdx.x-1)*2+1]); /*busy wait*/ 
+				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);
+				transblock_buffers[(blockIdx.x-1)*_2XBUFFER_SIZE+BUFFER_SIZE+threadIdx.x] = left_outgoing[threadIdx.x];
+				atomicExch(buffer_flags + (blockIdx.x-1) * 2 + 1, 1);
+				atomicExch(buffer_flags + (blockIdx.x-1) * 2, 0);
 				break;
 			case RBUFF:
-				while (buffer_flags[blockIdx.x*2]); /*busy wait*/
+				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];


static __device__ int blocksComplete = 0;

extern "C" __global__ void ParticleSort (unsigned int *global_mem,
					 unsigned int *transblock_buffers,
					 unsigned int *buffer_flags,

	else if ((blockIdx.x < gridDim.x - 1) && (threadIdx.x >= (blockDim.x - BUFFER_SIZE)))
		buff_role = RBUFF;
	else buff_role = NONE;

	volatile unsigned int *const here = beginning + threadIdx.x;

