 					= right_outgoing[(cur_R_outgoing_written-TRANSFER_SIZE+1-blockDim.x+TRANSFER_SIZE+threadIdx.x)&BUFFER_MOD];
 				if (role == RIGHT)
 					atomicAdd(global_right_outgoing_counter, TRANSFER_SIZE);
-#ifdef DEBUG
-				if (iter < TEST_ITER && role == RIGHT) {
-					for (int i = 0; i < _2XBUFFER_SIZE; i++)
-						buffer_snapshot[iter][i] = global_right_outgoing_buffer[i];
-					++iter;
-				}
-#endif
 				if (*global_right_incoming_counter > cur_R_incoming_written) {
 					right_incoming[(cur_R_incoming_written-blockDim.x+TRANSFER_SIZE+threadIdx.x)&BUFFER_MOD]
 						= global_right_incoming_buffer[(*global_right_incoming_counter-TRANSFER_SIZE+1+threadIdx.x)&BUFFER_MOD];


static __device__ int blocksComplete = 0;

#ifdef DEBUG
#define TEST_ITER 3
	__device__ unsigned int buffer_snapshot [TEST_ITER][_2XBUFFER_SIZE];
	__device__ int iter = 0; 
	__device__ int buffer_read = 0;
#endif

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

				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];
				if (role == LEFT) 
					atomicAdd(global_left_outgoing_counter, TRANSFER_SIZE);
				if (*global_left_incoming_counter > cur_L_incoming_written) {
