 			"cudaMemcpy device->host");
 #ifdef DEBUG
 	unsigned int *buff_ct = (unsigned int *)malloc(buffer_flags_size);
+	int buffer_read;
+
 	cudaMemcpy(buff_ct, buffer_flags, buffer_flags_size, cudaMemcpyDeviceToHost);
 	fprintf(stderr, "Buffer 1 ct A: %u, ct B: %u\n", buff_ct[0], buff_ct[1]);
+	
+	cudaMemcpyFromSymbol(&buffer_read, "buffer_read", sizeof(int), 0, cudaMemcpyDeviceToHost);
+	fprintf(stderr, "Number of buffer reads: %i\n", buffer_read);
 
 	free(buff_ct);
+
+	unsigned int buffer_snapshot[TEST_ITER][_2XBUFFER_SIZE];
+	cudaMemcpyFromSymbol(buffer_snapshot, "buffer_snapshot", TEST_ITER * _2XBUFFER_SIZE * sizeof(int), 0, cudaMemcpyDeviceToHost);
+
+	for (int i = 0; i < TEST_ITER; i++) {
+		fprintf(stderr, "Step %i:\n[[", i);
+		for (int j = 0; j < BUFFER_SIZE; j++)
+			fprintf(stderr, "%u, ", buffer_snapshot[i][j] & COLOR_MASK);
+		fprintf(stderr, "]\n [");
+		for (int j = BUFFER_SIZE; j < _2XBUFFER_SIZE; j++)
+			fprintf(stderr, "%u, ", buffer_snapshot[i][j] & COLOR_MASK);
+		fprintf(stderr, "]\n\n");
+	}
 #endif
 
 	ErrorCheck(cudaFree(global_mem), "cudaFree global");

				break;
			case RBUFF:
				global_right_outgoing_buffer[(*global_right_outgoing_counter-blockDim.x+TRANSFER_SIZE+threadIdx.x)&BUFFER_MOD]
					= left_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);
				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];

			case RIGHT:
				if (cur_R_incoming_read < cur_R_incoming_written)
					*here = right_incoming[cur_R_incoming_read++ & BUFFER_MOD];
#ifdef DEBUG
				atomicAdd(&buffer_read, 1);
#endif
				break;
			case END:
				if (going_right.color)
