 	
 __device__ inline bool operator< (const bin_id_pair& a, const bin_id_pair& b)
 	{
-	return (a.bin < b.bin);
+	if (a.bin == b.bin)
+		return (a.id < b.id);
+	else
+		return (a.bin < b.bin);
 	}
 
 __device__ inline bool operator> (const bin_id_pair& a, const bin_id_pair& b)
 	{
-	return (a.bin > b.bin);
+	if (a.bin == b.bin)
+		return (a.id > b.id);
+	else
+		return (a.bin > b.bin);	
+	}
+
+template<class T, unsigned int block_size> __device__ inline void scan_naive(T *temp)
+	{
+	int thid = threadIdx.x;
+	int n = blockDim.x;
+	
+	int pout = 0;
+	int pin = 1;
+	
+	for (int offset = 1; offset < block_size; offset *= 2)
+		{
+		pout = 1 - pout;
+		pin  = 1 - pout;
+		__syncthreads();
+		
+		temp[pout*n+thid] = temp[pin*n+thid];
+		
+		if (thid >= offset)
+			temp[pout*n+thid] += temp[pin*n+thid - offset];
+		}
+		
+	__syncthreads();
+	// bring the data back to the initial array
+	if (pout == 1)
+		{
+		pout = 1 - pout;
+		pin  = 1 - pout;
+		temp[pout*n+thid] = temp[pin*n+thid];
+		__syncthreads();
+		}
 	}
 
 /*template<unsigned int block_size> */__global__ void rebin_simple_sort_kernel(unsigned int *d_idxlist, unsigned int *d_bin_size, float4 *d_pos, unsigned int N, float xlo, float ylo, float zlo, unsigned int Mx, unsigned int My, unsigned int Mz, unsigned int Nmax, float scalex, float scaley, float scalez)
 	{
-	const int block_size = 256;
+	const int block_size = 32;
 	// read in the particle that belongs to this thread
 	unsigned int idx = blockDim.x * blockIdx.x + threadIdx.x;
 	

	initialize_data();
	sort_data();
	
	// run the various benchmarks
	bmark_host_rebinning(false);
	bmark_host_rebinning(true);

	// verify results again to be sure
	if (!verify())
		{
		printf("Invalid results at end of GPU/simple/sort bmark!\n");
		return;
		}	
	
