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

			printf("%d %d\n", sdata[i].bin, sdata[i].id);
			}
		}*/
	
	unsigned int size = atomicInc(&d_bin_size[bin], 0xffffffff);
	if (size < Nmax)
		d_idxlist[bin*Nmax + size] = idx;
	}
	
void rebin_particles_simple_sort(unsigned int *idxlist, unsigned int *bin_size, float4 *pos, unsigned int N, float Lx, float Ly, float Lz, unsigned int Mx, unsigned int My, unsigned int Mz, unsigned int Nmax)
	{
	// run one particle per thread
	const int block_size = 256;
	int n_blocks = (int)ceil(float(N)/(float)block_size);

	// make even bin dimensions

	
__device__ inline bool operator< (const bin_id_pair& a, const bin_id_pair& b)
	{
	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)
	{
	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 = 32;
	// read in the particle that belongs to this thread
	unsigned int idx = blockDim.x * blockIdx.x + threadIdx.x;
	
