 #define BLOCK 416
 #define BUFFER_SIZE 256
 #define TRANSFER_SIZE 16
-#define BUFFER_MOD 255
+#define BUFFER_MOD 0xFF
 #define _2XBUFFER_SIZE (BUFFER_SIZE * 2)
 #define MAX_MOMENTUM 0xF
 #define MOMENTUM_INIT 0xF0000000

static __device__ void Swap (struct particle *, struct particle *);

static __device__ int blocksComplete = 0;

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

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