 
 //! Texture for reading particle positions
 texture<float4, 1, cudaReadModeElementType> pdata_pos_tex;
-
+__constant__ float* electronDensity;
+__constant__	float* pairPotential;
+__constant__	float* embeddingFunction;
+__constant__	float* derivativeElectronDensity;
+__constant__	float* derivativePairPotential;
+__constant__	float* derivativeEmbeddingFunction;
+__constant__	float* atomDerivativeEmbeddingFunction;
+__constant__	EAMData eam_data;
 //!First kernel need to calculate a part of pair potential and to calculate
 /*!  Derivative Embedding Function for each atom.
 */

	cudaError_t error = cudaBindTexture(0, pdata_pos_tex, pdata.pos, sizeof(float4) * pdata.N);
	if (error != cudaSuccess)
		return error;

    gpu_compute_eam_forces_kernel<<< grid, threads, sizeof(float2)*coeff_width*coeff_width >>>(force_data,
	pdata,
	box,
	nlist,
	d_coeffs,
	eam_arrays.electronDensity,
	eam_arrays.pairPotential,
	eam_arrays.embeddingFunction,
	eam_arrays.derivativeElectronDensity,
	eam_arrays.derivativePairPotential,
	eam_arrays.derivativeEmbeddingFunction,
	eam_arrays.atomDerivativeEmbeddingFunction,
	eam_data);

	cudaThreadSynchronize();


	gpu_pdata_arrays pdata,
	gpu_boxsize box,
	gpu_nlist_array nlist,
	float2 *d_coeffs)
	{
	// read in the coefficients
	extern __shared__ float2 s_coeffs[];
