 					float dz = pz - bin_sdata[cur_offset + 2*blockDim.x];
 					dz = dz - box.Lz * rintf(dz * box.Lzinv);
 
-					int cur_neigh = __float_as_int(bin_sdata[cur_offset + 3*blockDim.x]);
-					
 					float dr = dx*dx + dy*dy + dz*dz;
 					
-					if ((my_pidx != EMPTY_BIN) && dr < r_maxsq && (cur_neigh != EMPTY_BIN) && (my_pidx != cur_neigh))
+					int not_excluded = (exclude.x != cur_neigh) & (exclude.y != cur_neigh) & (exclude.z != cur_neigh) & (exclude.w != cur_neigh);
+					if (dr < r_maxsq && (my_pidx != cur_neigh) && not_excluded)
 						{
 						nlist.list[my_pidx + (1 + n_neigh)*nlist.pitch] = cur_neigh;
 						n_neigh++;
 						}
 					}
+					}
 				}
 			}
 		}
-	// now, at the very end, we need to write out the number of particles in each list
-	nlist.list[my_pidx] = n_neigh;
-	nlist.last_updated_pos[my_pidx] = pdata.pos[my_pidx];
-	}*/
+		
+	if (my_pidx != EMPTY_BIN)
+		{
+		// now, at the very end, we need to write out the number of particles in each list
+		nlist.list[my_pidx] = n_neigh;
+		nlist.last_updated_pos[my_pidx] = pdata.pos[my_pidx];
+		}
+	}
 
 // non-unrolled version (with idx full) 
 /*__global__ void updateFromBins(gpu_pdata_arrays pdata, gpu_bin_array bins, gpu_nlist_array nlist, float r_maxsq, gpu_boxsize box)

// the shared data will hold the particle positions of all the particles in the neighboring bins as well as the 
// particle index of the particle (stored as an int to be accessed with __float_as_int)
// px,py,pz will have to be checked vs all of these
extern __shared__ float bin_sdata[];

__device__ int wrap_bin_i(int i, int M)
	{
	if (i < 0)
		i += M;

				int neigh_pidx = bins.idxlist[neigh_bin*bins.Nmax + threadIdx.x];
				
				// read in the neighbor position
				float4 neigh_pos = tex1Dfetch(pdata_pos_tex, neigh_pidx);
				// put it into shared memory
				
