codekingpro/portable-devtools

1/******************************************************************************2 * Copyright (c) 2020, NVIDIA CORPORATION.  All rights reserved.3 *4 * Redistribution and use in source and binary forms, with or without5 * modification, are permitted provided that the following conditions are met:6 *     * Redistributions of source code must retain the above copyright7 *       notice, this list of conditions and the following disclaimer.8 *     * Redistributions in binary form must reproduce the above copyright9 *       notice, this list of conditions and the following disclaimer in the10 *       documentation and/or other materials provided with the distribution.11 *     * Neither the name of the NVIDIA CORPORATION nor the12 *       names of its contributors may be used to endorse or promote products13 *       derived from this software without specific prior written permission.14 *15 * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"16 *AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE17 *IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE18 * DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY19 * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES20 * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;21 * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND22 * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT23 * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS24 * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.25 *26 ******************************************************************************/27 28/**29 * \file30 * Define helper math functions.31 */32 33#pragma once34 35#include <cub/config.cuh>36 37#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC)38#  pragma GCC system_header39#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG)40#  pragma clang system_header41#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC)42#  pragma system_header43#endif // no system header44 45#include <type_traits>46 47CUB_NAMESPACE_BEGIN48 49namespace detail50{51 52template <typename T>53using is_integral_or_enum =54  std::integral_constant<bool,55                         std::is_integral<T>::value || std::is_enum<T>::value>;56 57__host__ __device__ __forceinline__ constexpr  std::size_t58VshmemSize(std::size_t max_shmem,59           std::size_t shmem_per_block,60           std::size_t num_blocks)61{62  return shmem_per_block > max_shmem ? shmem_per_block * num_blocks : 0;63}64 65}66 67/**68 * Divide n by d, round up if any remainder, and return the result.69 *70 * Effectively performs `(n + d - 1) / d`, but is robust against the case where71 * `(n + d - 1)` would overflow.72 */73template <typename NumeratorT, typename DenominatorT>74__host__ __device__ __forceinline__ constexpr NumeratorT75DivideAndRoundUp(NumeratorT n, DenominatorT d)76{77  static_assert(cub::detail::is_integral_or_enum<NumeratorT>::value &&78                cub::detail::is_integral_or_enum<DenominatorT>::value,79                "DivideAndRoundUp is only intended for integral types.");80 81  // Static cast to undo integral promotion.82  return static_cast<NumeratorT>(n / d + (n % d != 0 ? 1 : 0));83}84 85constexpr __device__ __host__ int86Nominal4BItemsToItemsCombined(int nominal_4b_items_per_thread, int combined_bytes)87{88  return (cub::min)(nominal_4b_items_per_thread,89                    (cub::max)(1,90                               nominal_4b_items_per_thread * 8 /91                               combined_bytes));92}93 94template <typename T>95constexpr __device__ __host__ int96Nominal4BItemsToItems(int nominal_4b_items_per_thread)97{98  return (cub::min)(nominal_4b_items_per_thread,99                    (cub::max)(1,100                               nominal_4b_items_per_thread * 4 /101                                 static_cast<int>(sizeof(T))));102}103 104template <typename ItemT>105constexpr __device__ __host__ int106Nominal8BItemsToItems(int nominal_8b_items_per_thread)107{108  return sizeof(ItemT) <= 8u109           ? nominal_8b_items_per_thread110           : (cub::min)(nominal_8b_items_per_thread,111                        (cub::max)(1,112                                   ((nominal_8b_items_per_thread * 8) +113                                    static_cast<int>(sizeof(ItemT)) - 1) /114                                     static_cast<int>(sizeof(ItemT))));115}116 117/**118 * \brief Computes the midpoint of the integers119 *120 * Extra operation is performed in order to prevent overflow.121 *122 * \return Half the sum of \p begin and \p end123 */124template <typename T>125constexpr __device__ __host__ T MidPoint(T begin, T end)126{127  return begin + (end - begin) / 2;128}129 130CUB_NAMESPACE_END131