 
 	float4 my_pos = make_float4(0.0f, 0.0f, 0.0f, 0.0f);
 	uint4 exclude = make_uint4(EMPTY_BIN, EMPTY_BIN, EMPTY_BIN, EMPTY_BIN);
-	if (my_pidx != EMPTY_BIN)
-		{
+	//if (my_pidx != EMPTY_BIN)
+		//{
 		my_pos = tex1Dfetch(pdata_pos_tex, my_pidx);
 		//my_pos = pdata.pos[my_pidx];
 	
 		exclude = tex1Dfetch(nlist_exclude_tex, my_pidx);
 		//exclude = nlist.exclusions[my_pidx];
-		}
+		//}
 
 	float px = my_pos.x;
 	float py = my_pos.y;

	
	// read in the neighbor position
	float4 neigh_pos = make_float4(0.0f,0.0f,0.0f,0.0f);
	if (neigh_pidx != EMPTY_BIN)
		{
		neigh_pos = tex1Dfetch(pdata_pos_tex, neigh_pidx);
		//neigh_pos = pdata.pos[neigh_pidx];
		}
	
	// put it into shared memory
	
	__syncthreads();
	
	bin_sdata[threadIdx.x] = neigh_pos.x;

					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)
