         d_n_keep,
         nrecv);
     alloc.deallocate((char *) d_temp_storage);
- 
+
     hipMemcpy(&n_keep, d_n_keep, sizeof(unsigned int), hipMemcpyDeviceToHost);
     alloc.deallocate((char *)d_n_keep);
 

    const BoxDim& box)
    {
    assert(d_ghost_idx_adj);
    assert(d_tag);
    assert(d_pos);
    assert(d_img);
    assert(d_vel);
    assert(d_charge);
    assert(d_diameter);
    assert(d_body);
    assert(d_orientation);

    unsigned int block_size = 256;
    unsigned int n_blocks = n_out/block_size + 1;
    if (send_tag)
        {
        assert(d_tag_sendbuf);
        hipLaunchKernelGGL(gpu_pack_kernel, dim3(n_blocks), dim3(block_size), 0, 0,
            n_out, d_ghost_idx_adj, d_tag, d_tag_sendbuf);
        }
    if (send_pos)
        {
        assert(d_pos_sendbuf);
        hipLaunchKernelGGL(gpu_pack_wrap_kernel, dim3(n_blocks), dim3(block_size), 0, 0,
            n_out, d_ghost_idx_adj, d_pos, d_img, d_pos_sendbuf, send_image ? d_img_sendbuf : 0, di, my_pos, box);
        }
    if (send_vel)
        {
        assert(d_vel_sendbuf);
        hipLaunchKernelGGL(gpu_pack_kernel, dim3(n_blocks), dim3(block_size), 0, 0,
           n_out, d_ghost_idx_adj, d_vel, d_vel_sendbuf);
        }
    if (send_charge)
        {
        assert(d_charge_sendbuf);
        hipLaunchKernelGGL(gpu_pack_kernel, dim3(n_blocks), dim3(block_size), 0, 0,
            n_out, d_ghost_idx_adj, d_charge, d_charge_sendbuf);
        }
    if (send_diameter)
        {
        assert(d_diameter_sendbuf);
        hipLaunchKernelGGL(gpu_pack_kernel, dim3(n_blocks), dim3(block_size), 0, 0,
            n_out, d_ghost_idx_adj, d_diameter, d_diameter_sendbuf);
        }
    if (send_body)
        {
        assert(d_body_sendbuf);
        hipLaunchKernelGGL(gpu_pack_kernel, dim3(n_blocks), dim3(block_size), 0, 0,
            n_out, d_ghost_idx_adj, d_body, d_body_sendbuf);
        }
    if (send_orientation)
        {
        assert(d_orientation_sendbuf);
        hipLaunchKernelGGL(gpu_pack_kernel, dim3(n_blocks), dim3(block_size), 0, 0,
            n_out, d_ghost_idx_adj, d_orientation, d_orientation_sendbuf);

        d_n_out,
        n_groups);
    alloc.deallocate((char *) d_temp_storage);

    hipMemcpy(&n_out, d_n_out, sizeof(unsigned int), hipMemcpyDeviceToHost);
    alloc.deallocate((char *)d_n_out);
    }
