 	// but reading it from the texture gives a slightly better performance, possibly because it "warms up" the
 	// texture cache for the next read
 	float4 pos = make_float4(0.0f, 0.0f, 0.0f, 0.0f);
-	if (pidx < pdata.N)
-		{
-		#ifdef USE_CUDA_BUG_WORKAROUND
-		pos = pdata.pos[pidx];
-		#else
-		pos = tex1Dfetch(pdata_pos_tex, pidx);
-		#endif
-		}
-
+	pos = tex1Dfetch(pdata_pos_tex, pidx);
+	
 	// initialize the force to 0
 	float fx = 0.0f;
 	float fy = 0.0f;

	return retval;
	}

//! Hardcode the warp size for the divergence avoiding code
#define WARP_SIZE 32

//! Kernel for calculating lj forces
/*! This kerenel is called to calculate the lennard-jones forces on all N particles


	This version dynamically avoids divergent warps
	It is also assumed that the neighborlist is padded out to have a number of columsn a 
	multiple of the block size so that reads can be simply done in a big block
*/
extern "C" __global__ void calcLJForces_kernel(float4 *d_forces, gpu_pdata_arrays pdata, gpu_nlist_array nlist, float r_cutsq, gpu_boxsize box)
	{
	// start by identifying which particle we are to handle
	int pidx = blockIdx.x * blockDim.x + threadIdx.x;
	
	if (pidx >= pdata.N)
		return;
	
	// load in the length of the list (each thread loads it individually, 
	// we are going to dynamically avoid divergent warps by finding the maximum
