     cufftExecC2C(es_data.plan, es_data.GPU_rho_real_space, es_data.GPU_rho_real_space, CUFFT_FORWARD);
     cudaThreadSynchronize();
     
-    combined_green_e_kernel <<< new_gridsize, new_blockzise >>> (es_data.GPU_E_x, es_data.GPU_E_y, es_data.GPU_E_z, es_data.GPU_k_vec, es_data.GPU_rho_real_space,  es_data.Nx, es_data.Ny, es_data.Nz, es_data.GPU_green_hat);
+    combined_green_e_kernel <<< new_gridsize, new_blocksize >>> (es_data.GPU_E_x, es_data.GPU_E_y, es_data.GPU_E_z, es_data.GPU_k_vec, es_data.GPU_rho_real_space,  es_data.Nx, es_data.Ny, es_data.Nz, es_data.GPU_green_hat);
        
 
     //backtransform field:

void electrostatics_calculation(const gpu_force_data_arrays& force_data, const gpu_pdata_arrays &pdata, const gpu_boxsize &box, const float3 *d_params, const float *d_rcutsq)
{
    //first time allocation of memory-------------------------------------
    int blocksize = 512;
    dim3 grid( pdata.N / blocksize + 1, 1, 1);
    dim3 threads(blocksize, 1, 1);


	
	CPU_virial_energy.y -= CPU_virial_energy.z * es_data.kappa / sqrt(M_PI);

	set_to_zero <<< new_gridsize , new_blocksize >>> (es_data.GPU_rho_real_space, es_data.Nx, es_data.Ny, es_data.Nz);
	cudaThreadSynchronize();  
	return CPU_virial_energy;
	
