 				WriteParticle(&going_left, here - 1);
 			}
 		}
-		if ((role != IDLE) && !resident)
+		if ((role != IDLE) && !resident) {
 			isNotComplete = TRUE;
-		__syncthreads();
+			if (wasComplete) {
+				wasComplete = FALSE;
+				atomicAdd(&blocksComplete, -1);
+			}
+		}
 		if (i & 0xF == 0) {
 			switch (buff_role) {
 			case LBUFF:

	__shared__ unsigned int right_incoming [BUFFER_SIZE];
	__shared__ unsigned int right_outgoing [BUFFER_SIZE];
	__shared__ unsigned int isNotComplete;


	/* define registers */

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,
