     }
 
 
+//! Kernel to select ghost atoms due to non-bonded interactions
+__global__ void gpu_nonbonded_plan_kernel(unsigned char *plan,
+                               const unsigned int N,
+                               Scalar4 *d_postype,
+                               const BoxDim box,
+                               Scalar3 ghost_fraction)
+    {
+    unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;
+
+    if (idx >= N) return;
+
+    Scalar4 postype = d_postype[idx]; 
+    Scalar3 pos = make_scalar3(postype.x,postype.y,postype.z);
+    Scalar3 f = box.makeFraction(pos);
+
+    unsigned char p = plan[idx];
+
+    // is particle inside ghost layer? set plan accordingly.
+    if (f.x >= Scalar(1.0) - ghost_fraction.x)
+        p |= send_east;
+    if (f.x < ghost_fraction.x)
+        p |= send_west;
+    if (f.y >= Scalar(1.0) - ghost_fraction.y)
+        p |= send_north;
+    if (f.y < ghost_fraction.y)
+        p |= send_south;
+    if (f.z >= Scalar(1.0) - ghost_fraction.z)
+        p |= send_up;
+    if (f.z < ghost_fraction.z)
+        p |= send_down;
+
+    // write out plan
+    plan[idx] = p;
+    }
+
 //! Construct plans for sending non-bonded ghost particles
 /*! \param d_plan Array of ghost particle plans
  * \param N number of particles to check

#include "BondData.cuh"

#include <thrust/device_ptr.h>
#include <thrust/iterator/permutation_iterator.h>
#include <thrust/scatter.h>
#include <thrust/iterator/counting_iterator.h>
#include <thrust/iterator/constant_iterator.h>
#include <thrust/iterator/zip_iterator.h>

using namespace thrust;


                                      const BoxDim &box,
                                      float r_ghost)
    {
    unsigned int block_size = 512;

    Scalar3 ghost_fraction = r_ghost/box.getL();
    gpu_nonbonded_plan_kernel<<<N/block_size+1, block_size>>>(d_plan,
                                                              N,
                                                              d_pos,
                                                              box,
                                                              ghost_fraction);
    }

//! Kernel to pack local particle data into ghost send buffers
