                                   gpu_prefactRef[threadID]);
 }
 
-__global__ void BoxReciprocalSetupGPU(double *gpu_x,
-                                      double *gpu_y,
-                                      double *gpu_z,
-                                      double *gpu_kx,
-                                      double *gpu_ky,
-                                      double *gpu_kz,
-                                      double atomNumber,
-                                      double *gpu_particleCharge,
-                                      double *gpu_sumRnew,
-                                      double *gpu_sumInew,
-                                      int imageSize)
-{
-  int threadID = blockIdx.x * blockDim.x + threadIdx.x;
-  if(threadID >= imageSize)
-    return;
-  int i;
-  double dotP;
-
-  gpu_sumRnew[threadID] = 0.0;
-  gpu_sumInew[threadID] = 0.0;
-  for(i = 0; i < atomNumber; i++) {
-    dotP = DotProductGPU(gpu_kx[threadID], gpu_ky[threadID], gpu_kz[threadID],
-                         gpu_x[i], gpu_y[i], gpu_z[i]);
-    gpu_sumRnew[threadID] += gpu_particleCharge[i] * cos(dotP);
-    gpu_sumInew[threadID] += gpu_particleCharge[i] * sin(dotP);
-  }
-}
-
-__global__ void BoxReciprocalGPU(double *gpu_prefact,
-                                 double *gpu_sumRnew,
-                                 double *gpu_sumInew,
-                                 double *gpu_energyRecip,
-                                 int imageSize)
-{
-  int threadID = blockIdx.x * blockDim.x + threadIdx.x;
-  if(threadID >= imageSize)
-    return;
-
-  gpu_energyRecip[threadID] = ((gpu_sumRnew[threadID] * gpu_sumRnew[threadID] +
-                                gpu_sumInew[threadID] * gpu_sumInew[threadID]) *
-                               gpu_prefact[threadID]);
-}
-
 #endif

  CUFREE(gpu_energyRecip);
  CUFREE(gpu_final_energyRecip);
  CUFREE(d_temp_storage);
}

void CallMolReciprocalGPU(VariablesCUDA *vars,

using namespace cub;

#define IMAGES_PER_BLOCK 64
#define PARTICLE_PER_BLOCK 64

#define FULL_MASK 0xffffffff

