 	// pull in the particle of this thread that will be compared to everything else
 	int my_pidx = bins.idxlist[my_bin*bins.Nmax + threadIdx.x];
 	
-	//float4 my_pos = tex2D(pdata_pos_tex, my_pidx & TEX_MASK, my_pidx >> LOG2_TEX_WIDTH);
 	float4 my_pos = tex1Dfetch(pdata_pos_tex, my_pidx);
 	float px = my_pos.x;
 	float py = my_pos.y;
 	float pz = my_pos.z;
+	
+	uint4 exclude = tex1Dfetch(nlist_exclude_tex, my_pidx);
 
 	// 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)

					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))
						{
						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];
	}*/

// 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)

		nlist.list[my_pidx] = n_neigh;
		nlist.last_updated_pos[my_pidx] = pdata.pos[my_pidx];
		}
	}*/


// non-unrolled version 
extern "C" __global__ void updateFromBins(gpu_pdata_arrays pdata, gpu_bin_array bins, gpu_nlist_array nlist, float r_maxsq, gpu_boxsize box)
	{
	// each block is going to compute the neighborlist for all of the particles in blockIdx.x
	int my_bin = blockIdx.x;
