     {
     unsigned int idx = blockDim.x*blockIdx.x+threadIdx.x;
 
+    if (idx >= n_recv) return;
+
     if (send_tag) d_tag[idx] = d_tag_recvbuf[idx];
     if (send_pos) d_pos[idx] = d_pos_recvbuf[idx];
     if (send_vel) d_vel[idx] = d_vel_recvbuf[idx];

    bool send_orientation)
    {
    unsigned int buf_idx = blockIdx.x * blockDim.x + threadIdx.x;

    unsigned int tag = d_ghost_tag[buf_idx];
    unsigned int idx = d_rtag[tag];

    mgpu::ContextPtr mgpu_context,
    cached_allocator& alloc)
    {
    // wrap pointers
    thrust::device_ptr<const unsigned int> counts_ptr(d_counts);

    // temporary array for output neighbor ranks
    thrust::device_ptr<unsigned int> out_neighbor_ptr((unsigned int *)alloc.allocate(n_out*sizeof(unsigned int)));


    /*
     * expand each tag by the number of neighbors to send the corresponding ptl to
     * and assign each copy to a different neighbor
     */

    // allocate temporary array
    thrust::device_ptr<unsigned int> offsets_ptr((unsigned int *)alloc.allocate(N*sizeof(unsigned int)));

    // scan over counts
    thrust::exclusive_scan(thrust::cuda::par(alloc), counts_ptr, counts_ptr + N, offsets_ptr);

    gpu_expand_neighbors(n_out,
        thrust::raw_pointer_cast(offsets_ptr),
        d_tag, d_ghost_plan, N, d_ghost_tag,
        d_neighbors, d_adj, nneigh,
        thrust::raw_pointer_cast(out_neighbor_ptr),
        *mgpu_context);

    // sort tags by neighbors
    mgpu::LocalitySortPairs(thrust::raw_pointer_cast(out_neighbor_ptr),
                   d_ghost_tag, n_out, *mgpu_context);

    alloc.deallocate((char *)thrust::raw_pointer_cast(offsets_ptr),0);

    mgpu::SortedSearch<mgpu::MgpuBoundsLower>(d_unique_neighbors, n_unique_neigh,
        thrust::raw_pointer_cast(out_neighbor_ptr), n_out, d_ghost_begin, *mgpu_context);
    mgpu::SortedSearch<mgpu::MgpuBoundsUpper>(d_unique_neighbors, n_unique_neigh,
        thrust::raw_pointer_cast(out_neighbor_ptr), n_out, d_ghost_end, *mgpu_context);

    // deallocate temporary arrays
    alloc.deallocate((char *)thrust::raw_pointer_cast(out_neighbor_ptr),0);
    }
