 	// sort it 
 	bitonic_sort<bin_id_pair, block_size>(bin_pairs);
 	
-	// testing: print out the sorted data
-	/*if (idx == 0)
-		{
-		for (int i = 0; i < block_size; i++)
-			{
-			printf("%d %d\n", sdata[i].bin, sdata[i].id);
-			}
-		}*/
 	// identify the breaking points
 	__shared__ unsigned int unique[block_size*2];
 	

	// call the kernel
	//cudaMemset(gd_bin_size, 0, sizeof(unsigned int)*Mx*My*Mz);
	fast_memclear_kernal<<<(int)ceil(float(Mx*My*Mz)/(float)block_size), block_size>>>(gd_bin_size, Mx*My*Mz);
	rebin_simple_sort_kernel/*<block_size>*/<<<n_blocks, block_size>>>(idxlist, bin_size, pos, N, xlo, ylo, zlo, Mx, My, Mz, Nmax, scalex, scaley, scalez);
	}
	
// benchmark the device rebinning

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 = 64;
	int n_blocks = (int)ceil(float(N)/(float)block_size);

	// make even bin dimensions
