 	/* 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;
 		}

			// 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));

	ParticleSort<<<grid, block>>>(global_mem, transblock_buffers, buffer_flags, size);
	ErrorCheck(cudaMemcpy(buffer, global_mem, global_mem_size, cudaMemcpyDeviceToHost),
			"cudaMemcpy device->host");

	ErrorCheck(cudaFree(global_mem), "cudaFree global");
	ErrorCheck(cudaFree(transblock_buffers), "cudaFree buffers");
