 		bin = 0xffffffff;
 	
 	// load up shared memory
-	__shared__ bin_id_pair sdata[block_size];
-	sdata[threadIdx.x] = make_bin_id_pair(bin, idx);
+	__shared__ bin_id_pair bin_pairs[block_size];
+	bin_pairs[threadIdx.x] = make_bin_id_pair(bin, idx);
 	__syncthreads();
 	
 	// sort it 
-	bitonic_sort<bin_id_pair, block_size>(sdata);
+	bitonic_sort<bin_id_pair, block_size>(bin_pairs);
 	
 	// testing: print out the sorted data
 	/*if (idx == 0)

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

			printf("%d %d\n", sdata[i].bin, sdata[i].id);
			}
		}*/
	// identify the breaking points
	__shared__ unsigned int unique[block_size*2];
	
	bool is_unique = false;
	if (threadIdx.x > 0 && bin_pairs[threadIdx.x].bin != bin_pairs[threadIdx.x-1].bin)
		is_unique = true;
	
	unique[threadIdx.x] = 0;
	if (is_unique)
		unique[threadIdx.x] = 1;
	
	// threadIdx.x = 0 is unique: but we don't want to count it in the scan
	if (threadIdx.x == 0)
		is_unique = true;
	
	__syncthreads();
	
	// scan to find addresses to write to
	scan_naive<unsigned int, block_size>(unique);
	
	// determine start location of each unique value in the array
	__shared__ unsigned int start[block_size+1];
	
	if (is_unique)
		start[unique[threadIdx.x]] = threadIdx.x;
				
	// boundary condition: need one past the end
	if (threadIdx.x == 0)
		start[unique[block_size-1]+1] = block_size;
	
	__syncthreads();
	
	// now: each unique start point does it's own atomicAdd to find the starting offset
	__shared__ unsigned int start_offset[block_size];
	
	if (is_unique)
		start_offset[unique[threadIdx.x]] = atomicAdd(&d_bin_size[bin_pairs[threadIdx.x].bin], start[unique[threadIdx.x]+1] - start[unique[threadIdx.x]]);
	
	__syncthreads();
	
	// finally! we can write out all the particles
	unsigned int offset = start_offset[unique[threadIdx.x]];
	if (offset + threadIdx.x - start[unique[threadIdx.x]] < Nmax)
		d_idxlist[bin_pairs[threadIdx.x].bin*Nmax + offset + threadIdx.x - start[unique[threadIdx.x]]] = bin_pairs[threadIdx.x].id;
	}
	
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 = 32;
	int n_blocks = (int)ceil(float(N)/(float)block_size);

	// make even bin dimensions
