 	if (error != cudaSuccess)
 		return error;
     // run the kernel
-    cudaMemcpyToSymbol("eam_data", &eam_data, sizeof(EAMTexInterData));
+    cudaMemcpyToSymbol("eam_data_ti", &eam_data, sizeof(EAMTexInterData));
 
     gpu_compute_eam_tex_inter_forces_kernel<<< grid, threads>>>(force_data,
 	pdata,

//texture<float, 1, cudaReadModeElementType> derivativePairPotential_tex;
texture<float, 1, cudaReadModeElementType> derivativeEmbeddingFunction_tex;
texture<float, 1, cudaReadModeElementType> atomDerivativeEmbeddingFunction_tex;
__constant__ EAMTexInterData eam_data;

extern "C" __global__ void gpu_compute_eam_tex_inter_forces_kernel(
	gpu_force_data_arrays force_data,

		// calculate r squard (FLOPS: 5)
		float rsq = dx*dx + dy*dy + dz*dz;

		if (rsq > eam_data_ti.r_cutsq) continue;

		float inverseR = rsqrtf(rsq);
        float r = 1.0f / inverseR;
		position = r * eam_data_ti.rdr;
		int shift = (typei>=typej)?(int)((2 * ntypes - typej -1)*typej/2 + typei) * nr:(int)((2 * ntypes - typei -1)*typei/2 + typej) * nr;
        float2 pair_potential = tex1D(pairPotential_tex, position + shift + 0.5f);
		float pair_eng =  pair_potential.x * inverseR;

		float derivativePhi = (pair_potential.y - pair_eng) * inverseR;

		float derivativeRhoI = tex1D(derivativeElectronDensity_tex, position + typei * eam_data_ti.nr + 0.5f);

		float derivativeRhoJ = tex1D(derivativeElectronDensity_tex, position + typej * eam_data_ti.nr + 0.5f);

		float fullDerivativePhi = adef * derivativeRhoJ +
				atomDerivativeEmbeddingFunction[cur_neigh] * derivativeRhoI + derivativePhi;
