 	float* atomDerivativeEmbeddingFunction)
 	{
 	// start by identifying which particle we are to handle
-	int idx_local = blockIdx.x * blockDim.x + threadIdx.x;
-
-	if (idx_local >= pdata.local_num)
+	volatile  int idx = blockIdx.x * blockDim.x + threadIdx.x;
+
+	if (idx >= pdata.local_num)
 		return;
 
-	int idx_global = idx_local + pdata.local_beg;
-
 	// loadj in the length of the list (MEM_TRANSFER: 4 bytes)
-	int n_neigh = nlist.n_neigh[idx_global];
+	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_global);
+	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_global];
-	//float4 force = force_data.force[idx_local];
+	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_local].w;
+	//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_local].w;
+	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_global];
+	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_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);

	/*unsigned int r_index = (unsigned int)position;
	position -= (float)r_index;*/
	// . .
	atomDerivativeEmbeddingFunction[idx_global] = tex1D(derivativeEmbeddingFunction_tex, position + typei * eam_data.nrho + 0.5f);//derivativeEmbeddingFunction[r_index + typei * eam_data.nrho];

	force.w += tex1D(embeddingFunction_tex, position + typei * eam_data.nrho + 0.5f);//embeddingFunction[r_index + typei * eam_data.nrho] + derivativeEmbeddingFunction[r_index + typei * eam_data.nrho] * position * eam_data.drho;
	force_data.force[idx_local] = force;
	}
extern "C" __global__ void gpu_compute_eam_tex_inter_forces_kernel_2(
	gpu_force_data_arrays force_data,

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

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