     }
 
 
+//! 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

                                      const BoxDim &box,
                                      float r_ghost)
    {
    thrust::device_ptr<float4> pos_ptr(d_pos);
    thrust::device_ptr<unsigned char> plan_ptr(d_plan);

    thrust::transform(
        thrust::make_zip_iterator(thrust::make_tuple(
            pos_ptr,
            plan_ptr)),
        thrust::make_zip_iterator(thrust::make_tuple(
            pos_ptr,
            plan_ptr)) + N,
        plan_ptr,
        make_nonbonded_plan(box, r_ghost));
    }

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

    {
    return sizeof(update_element_gpu);
    }

unsigned int *d_n_fetch_ptl;              //!< Index of fetched particle from received ptl list

