       const auto tile_size = MergePolicyT::ITEMS_PER_TILE;
       const auto num_tiles = cub::DivideAndRoundUp(num_items, tile_size);
 
+      /**
+       * Merge sort supports large types, which can lead to excessive shared
+       * memory size requirements. In these cases, merge sort allocates virtual
+       * shared memory that resides in global memory:
+       * ```
+       * extern __shared__ char shmem[];
+       * typename AgentT::TempStorage &storage =
+       *   *reinterpret_cast<typename AgentT::TempStorage *>(
+       *     UseVShmem ? vshmem + vshmem_offset : shmem);
+       * ```
+       * Having `UseVShmem` as a runtime variable leads to the generation of
+       * generic loads and stores, which causes a slowdown. Therefore,
+       * `UseVShmem` has to be known at compilation time.
+       * In the generic case, available shared memory size is queried at runtime
+       * to check if kernels requirements are satisfied. Since the query result
+       * is not known at compile-time, merge sort kernels are specialized for
+       * both cases.
+       * To address increased compilation time, the dispatch layer checks
+       * whether kernels requirements fit into default shared memory
+       * size (48KB). In this case, there's no need for virtual shared
+       * memory specialization.
+       */
       constexpr std::size_t default_shared_memory_size = 48 * 1024;
       constexpr auto block_sort_shmem_size =
         static_cast<std::size_t>(BlockSortAgentT::SHARED_MEMORY_SIZE);

{
  using KeyT = typename std::iterator_traits<KeyIteratorT>::value_type;

  //------------------------------------------------------------------------------
  // Architecture-specific tuning policies
  //------------------------------------------------------------------------------

  struct Policy350 : ChainedPolicy<350, Policy350, Policy350>
  {

    }
  }

  // Constructor
  CUB_RUNTIME_FUNCTION __forceinline__
  DispatchMergeSort(void *d_temp_storage,
                    std::size_t &temp_storage_bytes,
