                                       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));
+    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

#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;


    {
    return sizeof(update_element_gpu);
    }

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

