 	__shared__ unsigned int right_outgoing [BUFFER_SIZE];
 	__shared__ unsigned int isNotComplete;
 	__shared__ unsigned int wasComplete;
-	__shared__ unsigned int *cur_left_incoming, *cur_left_outgoing;
-	__shared__ unsigned int *cur_right_incoming, *cur_right_outgoing;
+	__shared__ unsigned int cur_L_incoming_written, cur_L_incoming_read;
+	__shared__ unsigned int cur_L_outgoing_written, cur_L_outgoing_read;
+	__shared__ unsigned int cur_R_incoming_written, cur_R_incoming_read;
+	__shared__ unsigned int cur_R_outgoing_written, cur_R_outgoing_read;
+	__shared__ unsigned int global_L_buffer_read, global_R_buffer_read;
 
 
 	/* define registers */
 	const int absThreadID = blockIdx.x * blockDim.x + threadIdx.x;
 
-	/* these are only used by a couple of threads . . . */
 	if (threadIdx.x == 0) {
-		cur_left_incoming = left_incoming + BUFFER_SIZE - 1;
-		cur_left_outgoing = left_outgoing + BUFFER_SIZE - 1;
-		cur_right_incoming = right_incoming;
-		cur_right_outgoing = right_outgoing;
+		cur_L_incoming_written = cur_L_incoming_read = cur_L_outgoing_written = cur_L_outgoing_read =
+			cur_R_incoming_written = cur_R_incoming_read = cur_R_outgoing_written = cur_R_outgoing_read =
+			global_L_buffer_read = global_R_buffer_read = 0;
 	}
+	unsigned int *const global_left_incoming_buffer = transblock_buffers + (blockIdx.x-1) * _2XBUFFER_SIZE;
+	unsigned int *const global_left_outgoing_buffer = transblock_buffers + (blockIdx.x-1) * _2XBUFFER_SIZE + BUFFER_SIZE;
+	unsigned int *const global_right_outgoing_buffer = transblock_buffers + blockIdx.x * _2XBUFFER_SIZE;
+	unsigned int *const global_right_incoming_buffer = transblock_buffers + blockIdx.x * _2XBUFFER_SIZE + BUFFER_SIZE;
+	unsigned int *const global_left_incoming_counter = buffer_count + (blockIdx.x-1) * 2;
+	unsigned int *const global_left_outgoing_counter = buffer_count + (blockIdx.x-1) * 2 + 1;
+	unsigned int *const global_right_outgoing_counter = buffer_count + blockIdx.x * 2;
+	unsigned int *const global_right_incoming_counter = buffer_count + blockIdx.x * 2 + 1;
 	
 	struct particle going_left, going_right;
 


static __device__ int blocksComplete = 0;


extern "C" __global__ void ParticleSort (unsigned int *global_mem,
					 unsigned int *transblock_buffers,
					 unsigned int *buffer_flags,
					 unsigned long size)
{
	/* define shared memory */

	else role = MIDDLE;

	enum {LBUFF, NONE, RBUFF} buff_role;
	if ((blockIdx.x > 0) && (threadIdx.x < TRANSFER_SIZE))
		buff_role = LBUFF;
	else if ((blockIdx.x < gridDim.x - 1) && (threadIdx.x >= (blockDim.x - TRANSFER_SIZE)))
		buff_role = RBUFF;
	else buff_role = NONE;

