     unsigned int block_size = 256;
     unsigned int n_blocks = N/block_size + 1;
 
-    hipLaunchKernelGGL(HIP_KERNEL_NAME(gpu_select_particle_migrate), dim3(n_blocks), dim3(block_size), 0, 0, 
+    hipLaunchKernelGGL(HIP_KERNEL_NAME(gpu_select_particle_migrate), dim3(n_blocks), dim3(block_size), 0, 0,
         N,
         d_pos,
         d_comm_flag,

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

    unsigned int block_size = 256;
    unsigned int n_blocks = n_out/block_size + 1;
    if (send_tag)
        {
        assert(d_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);
        assert(d_pos_sendbuf);
        if (send_image)
            {
            assert(d_img);
            }
        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);
        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);
        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);
        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);
        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);
        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);
