 	/* initial coalesced global memory read */
 	if (role != IDLE) {
 		resident = MOMENTUM_INIT | (global_mem[absThreadID] + 1);
-		if (threadIdx.x & 0x01 || role == END) {
+		if (threadIdx.x & 0x01 || role == END || role == RIGHT) {
 			ReadParticle(resident, &going_left);
 			resident = 0;
 		}


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,

				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)
