 		// read the current neighbor index (MEM TRANSFER: 4 bytes)
 		// prefetch the next value and set the current one
 		cur_neigh = next_neigh;
-		next_neigh = nlist.list[nlist.pitch*(neigh_idx+1) + idx_global];
+		next_neigh = nlist.list[nlist.pitch*(neigh_idx+1) + idx];
 
 		// get the neighbor's position (MEM TRANSFER: 16 bytes)
 		float4 neigh_pos = tex1Dfetch(pdata_pos_tex, cur_neigh);

	force.z = fzi;
	force.w += m_pe;
	// now that the force calculation is complete, write out the result (MEM TRANSFER: 20 bytes)
	force_data.force[idx_local] = force;
	force_data.virial[idx_local] = virial;
	}

/*! \param force_data Force data on GPU to write forces to

	float* atomDerivativeEmbeddingFunction)
	{
	// start by identifying which particle we are to handle
	volatile  int idx = blockIdx.x * blockDim.x + threadIdx.x;

	if (idx >= pdata.local_num)
		return;

	// loadj in the length of the list (MEM_TRANSFER: 4 bytes)
	int n_neigh = nlist.n_neigh[idx];

	// read in the position of our particle. Texture reads of float4's are faster than global reads on compute 1.0 hardware
	// (MEM TRANSFER: 16 bytes)
	float4 pos = tex1Dfetch(pdata_pos_tex, idx);
	int typei = __float_as_int(pos.w);
	// prefetch neighbor index
	float position;
	int cur_neigh = 0;
	int next_neigh = nlist.list[idx];
	//float4 force = force_data.force[idx];
	float4 force = make_float4(0.0f, 0.0f, 0.0f, 0.0f);
	//force.w = force_data.force[idx].w;
	float fxi = 0.0f;
	float fyi = 0.0f;
	float fzi = 0.0f;
	float m_pe = 0.0f;
	float pairForce = 0.0f;
	float virial = 0.0f;
	force.w = force_data.force[idx].w;
	int nr = eam_data.nr;
	int nrho = eam_data.nrho;
	int ntypes = eam_data.ntypes;
	float adef = atomDerivativeEmbeddingFunction[idx];
	for (int neigh_idx = 0; neigh_idx < n_neigh; neigh_idx++)
		{
		cur_neigh = next_neigh;
		next_neigh = nlist.list[nlist.pitch*(neigh_idx+1) + idx];

		// get the neighbor's position (MEM TRANSFER: 16 bytes)
		float4 neigh_pos = tex1Dfetch(pdata_pos_tex, cur_neigh);
