 	else if ((blockIdx.x < gridDim.x - 1) && (threadIdx.x >= (blockDim.x - BUFFER_SIZE)))
 		buff_role = RBUFF;
 	else buff_role = NONE;
-
-	switch (buff_role) {
-	case LBUFF:
-		left_incoming[threadIdx.x] = 0;
-		left_outgoing[threadIdx.x] = 0;
-		break;
-	case RBUFF:
-		right_incoming[threadIdx.x-blockDim.x+BUFFER_SIZE] = 0;
-		right_outgoing[threadIdx.x-blockDim.x+BUFFER_SIZE] = 0;
-		break;
-	}
 
 	volatile unsigned int *const here = beginning + threadIdx.x;
 


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 + 1, 1);
				atomicExch(buffer_flags + (blockIdx.x-1) * 2, 0);
				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];
