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

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

					float dz = pz - bin_sdata[cur_offset + 2*blockDim.x];
					dz = dz - box.Lz * rintf(dz * box.Lzinv);

					float dr = dx*dx + dy*dy + dz*dz;
					
					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++;
						}
					}
					}
				}
			}
		}
		
	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)
