     assert(args.d_excell_size);
     assert(args.d_check_overlaps);
 
-    // determine the maximum block size and clamp the input block size down
-    static int max_block_size = -1;
-    static hipFuncAttributes attr;
-    if (max_block_size == -1)
-        {
-        hipFuncGetAttributes(&attr, reinterpret_cast<const void*>(&kernel::hpmc_insert_depletants<Shape>));
-        max_block_size = attr.maxThreadsPerBlock;
-        if (max_block_size % args.devprop.warpSize)
-            // handle non-sensical return values from hipFuncGetAttributes
-            max_block_size = (max_block_size/args.devprop.warpSize-1)*args.devprop.warpSize;
-        }
-
-    // choose a block size based on the max block size by regs (max_block_size) and include dynamic shared memory usage
-    unsigned int block_size = min(args.block_size, (unsigned int)max_block_size);
-
-    unsigned int tpp = min(args.tpp,block_size);
-    unsigned int n_groups = block_size / tpp;
-    unsigned int max_queue_size = n_groups*tpp;
-    unsigned int max_depletant_queue_size = n_groups;
-
-    const unsigned int min_shared_bytes = args.num_types * sizeof(typename Shape::param_type) +
-               args.overlap_idx.getNumElements() * sizeof(unsigned int);
-
-    unsigned int shared_bytes = n_groups *(sizeof(Scalar4) + sizeof(Scalar3)) +
-                                max_queue_size*2*sizeof(unsigned int) +
-                                max_depletant_queue_size*sizeof(unsigned int) +
-                                min_shared_bytes;
-
-    if (min_shared_bytes >= args.devprop.sharedMemPerBlock)
-        throw std::runtime_error("Insufficient shared memory for HPMC kernel: reduce number of particle types or size of shape parameters");
-
-    while (shared_bytes + attr.sharedSizeBytes >= args.devprop.sharedMemPerBlock)
-        {
-        block_size -= args.devprop.warpSize;
-        if (block_size == 0)
-            throw std::runtime_error("Insufficient shared memory for HPMC kernel");
-        tpp = min(tpp, block_size);
-        n_groups = block_size / tpp;
-        max_queue_size = n_groups*tpp;
-        max_depletant_queue_size = n_groups;
-
-        shared_bytes = n_groups * (sizeof(Scalar4) + sizeof(Scalar3)) +
-                       max_queue_size*2*sizeof(unsigned int) +
-                       max_depletant_queue_size*sizeof(unsigned int) +
-                       min_shared_bytes;
-        }
-
-    static unsigned int base_shared_bytes = UINT_MAX;
-    bool shared_bytes_changed = base_shared_bytes != shared_bytes + attr.sharedSizeBytes;
-    base_shared_bytes = shared_bytes + attr.sharedSizeBytes;
-
-    unsigned int max_extra_bytes = args.devprop.sharedMemPerBlock - base_shared_bytes;
-    static unsigned int extra_bytes = UINT_MAX;
-    if (extra_bytes == UINT_MAX || args.update_shape_param || shared_bytes_changed)
-        {
-        // required for memory coherency
-        hipDeviceSynchronize();
-
-        // determine dynamically requested shared memory
-        char *ptr = (char *) nullptr;
-        unsigned int available_bytes = max_extra_bytes;
-        for (unsigned int i = 0; i < args.num_types; ++i)
-            {
-            params[i].allocate_shared(ptr, available_bytes);
-            }
-        extra_bytes = max_extra_bytes - available_bytes;
-        }
-
-    shared_bytes += extra_bytes;
-
-    // setup the grid to run the kernel
-    dim3 threads(tpp, n_groups,1);
-
-    for (int idev = args.gpu_partition.getNumActiveGPUs() - 1; idev >= 0; --idev)
-        {
-        auto range = args.gpu_partition.getRangeAndSetGPU(idev);
-
-        if (range.first == range.second)
-            continue;
-
-        unsigned int blocks_per_particle = implicit_args.max_n_depletants[idev] /
-            (implicit_args.depletants_per_group*n_groups) + 1;
-
-        dim3 grid( range.second-range.first, blocks_per_particle, 1);
-
-        if (blocks_per_particle > args.devprop.maxGridSize[1])
-            {
-            grid.y = args.devprop.maxGridSize[1];
-            grid.z = blocks_per_particle/args.devprop.maxGridSize[1]+1;
-            }
-
-        hipLaunchKernelGGL((kernel::hpmc_insert_depletants<Shape>), dim3(grid), dim3(threads), shared_bytes, implicit_args.streams[idev],
-                                                                     args.d_trial_postype,
-                                                                     args.d_trial_orientation,
-                                                                     args.d_trial_move_type,
-                                                                     args.d_postype,
-                                                                     args.d_orientation,
-                                                                     args.d_counters,
-                                                                     args.d_excell_idx,
-                                                                     args.d_excell_size,
-                                                                     args.excli,
-                                                                     args.cell_dim,
-                                                                     args.ghost_width,
-                                                                     args.ci,
-                                                                     args.N + args.N_ghost,
-                                                                     args.N,
-                                                                     args.num_types,
-                                                                     args.seed,
-                                                                     args.d_check_overlaps,
-                                                                     args.overlap_idx,
-                                                                     args.timestep,
-                                                                     args.dim,
-                                                                     args.box,
-                                                                     args.select,
-                                                                     args.d_reject_out_of_cell,
-                                                                     params,
-                                                                     max_queue_size,
-                                                                     max_extra_bytes,
-                                                                     implicit_args.depletant_type_a,
-                                                                     implicit_args.depletant_type_b,
-                                                                     implicit_args.depletant_idx,
-                                                                     implicit_args.d_implicit_count + idev*implicit_args.implicit_counters_pitch,
-                                                                     implicit_args.d_lambda,
-                                                                     args.d_nneigh,
-                                                                     args.d_nlist,
-                                                                     args.maxn,
-                                                                     args.d_overflow,
-                                                                     implicit_args.repulsive,
-                                                                     range.first,
-                                                                     max_depletant_queue_size,
-                                                                     implicit_args.d_n_depletants);
-        }
+    // select the kernel template according to the next power of two of the block size
+    unsigned int launch_bounds = WARP_SIZE;
+    while (launch_bounds < args.block_size)
+        launch_bounds *= 2;
+
+    kernel::depletants_launcher<Shape>(args, implicit_args, params, launch_bounds, int2type<MAX_BLOCK_SIZE>());
     }
 
 //! Driver for kernel::hpmc_update_pdata()

    }
#endif

} // end namespace gpu

} // end namespace hpmc

namespace hpmc {

namespace gpu {

// currently this is hardcoded, we should set it to the max of platforms
#if defined(__HIP_PLATFORM_NVCC__)
#define MAX_BLOCK_SIZE 1024
#define WARP_SIZE 32
#elif defined(__HIP_PLATFORM_HCC__)
#define MAX_BLOCK_SIZE 1024
#define WARP_SIZE 64
#endif

template<unsigned int> struct int2type{};

//! Wraps arguments to hpmc_* template functions
/*! \ingroup hpmc_data_structs */
