 				atomicAdd(&blocksComplete, -1);
 			}
 		}
+		__syncthreads();
 		if ((++i & 0x1F) == 0) {
 			switch (buff_role) {
 			case LBUFF:		
-				global_left_outgoing_buffer[(*global_left_outgoing_counter + threadIdx.x) & BUFFER_MOD]
-					= left_outgoing[(cur_L_outgoing_written - TRANSFER_SIZE + 1 + threadIdx.x) & BUFFER_MOD];
+				global_left_outgoing_buffer[(*global_left_outgoing_counter+threadIdx.x)&BUFFER_MOD]
+					= left_outgoing[(cur_L_outgoing_written-TRANSFER_SIZE+1+threadIdx.x)&BUFFER_MOD];
 				if (role == LEFT) 
 					atomicAdd(global_left_outgoing_counter, TRANSFER_SIZE);
 				if (*global_left_incoming_counter > cur_L_incoming_written) {

			// prepare for moving right
			switch (role) {
			case LEFT:
				if (cur_L_incoming_read < cur_L_incoming_written)
					*here = left_incoming[cur_L_incoming_read++ & BUFFER_MOD];
				WriteParticle(&going_right, here + 1);
#ifdef DEBUG
				atomicAdd(&buffer_read, 1);
#endif
				break;
			case RIGHT:
				WriteParticle(&going_right, right_outgoing + (cur_R_outgoing_written++ & BUFFER_MOD));

				RESET(going_left);
				break;
			case RIGHT:
				WriteParticle(&going_left, here - 1);
				if (cur_R_incoming_read < cur_R_incoming_written) 
					*here = right_incoming[cur_R_incoming_read++ & BUFFER_MOD];
				else
					*here = 0;
				break;
			case END:
				if (going_right.color)
