include/corecel/Macros.hh

0001 //------------------------------- -*- C++ -*- -------------------------------//
0002 // Copyright Celeritas contributors: see top-level COPYRIGHT file for details
0003 // SPDX-License-Identifier: (Apache-2.0 OR MIT)
0004 //---------------------------------------------------------------------------//
0005 /*!
0006  * \file corecel/Macros.hh
0007  * \brief Language and compiler abstraction macro definitions.
0008  *
0009  * The Macros file defines cross-platform (CUDA, C++, HIP) macros that
0010  * expand to attributes depending on the compiler and build configuration.
0011  */
0012 //---------------------------------------------------------------------------//
0013 #pragma once
0014
0015 #include "corecel/Config.hh"
0016
0017 //---------------------------------------------------------------------------//
0018 //!@{
0019 //! \name Compiler type/version macros
0020
0021 /*!
0022  * \def CELER_FUNCTION
0023  *
0024  * Decorate a function that works on both host and device, with and without
0025  * NVCC. The name of this function and its siblings is based on the Kokkos
0026  * naming scheme.
0027  */
0028 #if defined(__NVCC__) || defined(__HIP__)
0029 #    define CELER_FUNCTION __host__ __device__
0030 #else
0031 #    define CELER_FUNCTION
0032 #endif
0033
0034 #if defined(__NVCC__)
0035 #    define CELER_FORCEINLINE __forceinline__
0036 #elif defined(_MSC_VER)
0037 #    define CELER_FORCEINLINE inline __forceinline
0038 #elif defined(__clang__) || defined(__GNUC__) || defined(__HIP__) \
0039     || defined(__INTEL_COMPILER)
0040 #    define CELER_FORCEINLINE inline __attribute__((always_inline))
0041 #else
0042 #    define CELER_FORCEINLINE inline
0043 #endif
0044
0045 //! Detection for the current compiler isn't supported yet
0046 #define CELER_COMPILER_UNKNOWN 0
0047 //! Compiling with clang, or a clang-based compiler defining __clang__ (hipcc)
0048 #define CELER_COMPILER_CLANG 1
0049 /*!
0050  * \def CELER_COMPILER
0051  *
0052  * Compare to on of the CELER_COMPILER_<compiler> macros to check
0053  * which compiler is in use.
0054  *
0055  * TODO: add and test more compilers as needed.
0056  */
0057 #if defined(__clang__)
0058 #    define CELER_COMPILER CELER_COMPILER_CLANG
0059 #else
0060 #    define CELER_COMPILER CELER_COMPILER_UNKNOWN
0061 #endif
0062
0063 /*!
0064  * \def CELER_FORCEINLINE_FUNCTION
0065  *
0066  * Like CELER_FUNCTION but forces inlining. Compiler optimizers usually can
0067  * tell what needs optimizing, but this function can provide speedups (and
0068  * smaller sampling profiles) when inlining optimizations are not enabled. It
0069  * should be used sparingly.
0070  */
0071 #define CELER_FORCEINLINE_FUNCTION CELER_FUNCTION CELER_FORCEINLINE
0072
0073 // For use in detail code only for thin wrapper functions
0074 #define CELER_FIF CELER_FORCEINLINE_FUNCTION
0075
0076 /*!
0077  * \def CELER_CONSTEXPR_FUNCTION
0078  *
0079  * Decorate a function that works on both host and device, with and without
0080  * NVCC, can be evaluated at compile time, and should be forcibly inlined.
0081  */
0082 #define CELER_CONSTEXPR_FUNCTION constexpr CELER_FORCEINLINE_FUNCTION
0083
0084 // For use in detail code only for thin wrapper functions
0085 #define CELER_CEF CELER_CONSTEXPR_FUNCTION
0086
0087 /*!
0088  * \def CELER_UNLIKELY(condition)
0089  *
0090  * Mark the result of this condition to be "unlikely".
0091  *
0092  * This asks the compiler to move the section of code to a "cold" part of the
0093  * instructions, improving instruction locality. It should be used primarily
0094  * for error checking conditions.
0095  */
0096 #if defined(__clang__) || defined(__GNUC__)
0097 // GCC and Clang support the same builtin
0098 #    define CELER_UNLIKELY(COND) __builtin_expect(!!(COND), 0)
0099 #else
0100 // No other compilers seem to have a similar builtin
0101 #    define CELER_UNLIKELY(COND) (COND)
0102 #endif
0103
0104 /*!
0105  * \def CELER_UNREACHABLE
0106  *
0107  * Mark a point in code as being impossible to reach in normal execution.
0108  *
0109  * See https://clang.llvm.org/docs/LanguageExtensions.html#builtin-unreachable
0110  * or https://msdn.microsoft.com/en-us/library/1b3fsfxw.aspx
0111  * or
0112  * https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#__builtin_unreachable
0113  *
0114  * (The "unreachable" and "assume" compiler optimizations for CUDA are only
0115  * available in API version 11.3 or higher, which is encoded as
0116  * \code major*1000 + minor*10 \endcode).
0117  *
0118  * \note This macro should not generally be used; instead, the macro \c
0119  * CELER_ASSERT_UNREACHABLE() defined in base/Assert.hh should be used instead
0120  * (to provide a more detailed error message in case the point *is* reached).
0121  */
0122 #if (!defined(__CUDA_ARCH__) && (defined(__clang__) || defined(__GNUC__))) \
0123     || defined(__NVCOMPILER)                                               \
0124     || (defined(__CUDA_ARCH__) && CUDART_VERSION >= 11030)                 \
0125     || defined(__HIP_DEVICE_COMPILE__)
0126 #    define CELER_UNREACHABLE __builtin_unreachable()
0127 #elif defined(_MSC_VER)
0128 #    define CELER_UNREACHABLE __assume(false)
0129 #else
0130 #    define CELER_UNREACHABLE
0131 #endif
0132
0133 /*!
0134  * \def CELER_USE_DEVICE
0135  *
0136  * True if HIP or CUDA are enabled, false otherwise.
0137  */
0138 #if CELERITAS_USE_CUDA || CELERITAS_USE_HIP
0139 #    define CELER_USE_DEVICE 1
0140 #else
0141 #    define CELER_USE_DEVICE 0
0142 #endif
0143
0144 /*!
0145  * \def CELER_DEVICE_SOURCE
0146  *
0147  * Defined and true if building a HIP or CUDA source file. This is a generic
0148  * replacement for \c __CUDACC__ .
0149  */
0150 #if defined(__CUDACC__) || defined(__HIP__)
0151 #    define CELER_DEVICE_SOURCE 1
0152 #elif defined(__DOXYGEN__)
0153 #    define CELER_DEVICE_SOURCE 0
0154 #endif
0155
0156 /*!
0157  * \def CELER_DEVICE_COMPILE
0158  *
0159  * Defined and true if building device code in HIP or CUDA. This is a generic
0160  * replacement for \c __CUDA_ARCH__ .
0161  */
0162 #if defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__)
0163 #    define CELER_DEVICE_COMPILE 1
0164 #elif defined(__DOXYGEN__)
0165 #    define CELER_DEVICE_COMPILE 0
0166 #endif
0167
0168 #if CELERITAS_USE_CUDA            \
0169     && (__CUDACC_VER_MAJOR__ < 11 \
0170         || (__CUDACC_VER_MAJOR__ == 11 && __CUDACC_VER_MINOR__ < 5))
0171 /*!
0172  * Work around older NVCC bugs with `if constexpr`.
0173  *
0174  * These cause errors such as \verbatim
0175  *    error: missing return statement at end of non-void function
0176  * \endverbatim
0177  */
0178 #    define CELER_CUDACC_BUGGY_IF_CONSTEXPR 1
0179 #else
0180 #    define CELER_CUDACC_BUGGY_IF_CONSTEXPR 0
0181 #endif
0182
0183 //!@}
0184 //---------------------------------------------------------------------------//
0185 //!@{
0186 //! \name Exception handling macros
0187
0188 /*!
0189  * \def CELER_TRY_HANDLE
0190  *
0191  * "Try" to execute the statement, and "handle" *all* thrown errors by calling
0192  * the given function-like error handler with a \c std::exception_ptr object.
0193  *
0194  * \note A file that uses this macro must include the \c \<exception\> header
0195  * (but since the \c HANDLE_EXCEPTION needs to take an exception pointer, it's
0196  * got to be included anyway).
0197  */
0198 #define CELER_TRY_HANDLE(STATEMENT, HANDLE_EXCEPTION)   \
0199     do                                                  \
0200     {                                                   \
0201         try                                             \
0202         {                                               \
0203             STATEMENT;                                  \
0204         }                                               \
0205         catch (...)                                     \
0206         {                                               \
0207             HANDLE_EXCEPTION(std::current_exception()); \
0208         }                                               \
0209     } while (0)
0210
0211 /*!
0212  * \def CELER_TRY_HANDLE_CONTEXT
0213  *
0214  * Try the given statement, and if it fails, chain it into the given exception.
0215  *
0216  * The given \c CONTEXT_EXCEPTION must be an expression that yields an rvalue
0217  * to a \c std::exception subclass that isn't \c final . The resulting chained
0218  * exception will be passed into \c HANDLE_EXCEPTION for processing.
0219  */
0220 #define CELER_TRY_HANDLE_CONTEXT(                          \
0221     STATEMENT, HANDLE_EXCEPTION, CONTEXT_EXCEPTION)        \
0222     CELER_TRY_HANDLE(                                      \
0223         do {                                               \
0224             try                                            \
0225             {                                              \
0226                 STATEMENT;                                 \
0227             }                                              \
0228             catch (...)                                    \
0229             {                                              \
0230                 std::throw_with_nested(CONTEXT_EXCEPTION); \
0231             }                                              \
0232         } while (0),                                       \
0233         HANDLE_EXCEPTION)
0234
0235 //!@}
0236 //---------------------------------------------------------------------------//
0237 //!@{
0238 //! \name Class definition macros
0239
0240 /*!
0241  * \def CELER_DEFAULT_COPY_MOVE
0242  *
0243  * Explicitly declare defaulted copy and move constructors and assignment
0244  * operators. Use this if the destructor is declared explicitly, or as part of
0245  * the "protected" section of an interface class to prevent assignment between
0246  * incompatible classes.
0247  */
0248 #define CELER_DEFAULT_COPY_MOVE(CLS)      \
0249     CLS(CLS const&) = default;            \
0250     CLS& operator=(CLS const&) = default; \
0251     CLS(CLS&&) = default;                 \
0252     CLS& operator=(CLS&&) = default
0253
0254 /*!
0255  * \def CELER_DELETE_COPY_MOVE
0256  *
0257  * Explicitly declare *deleted* copy and move constructors and assignment
0258  * operators. Use this for scoped RAII classes.
0259  */
0260 #define CELER_DELETE_COPY_MOVE(CLS)      \
0261     CLS(CLS const&) = delete;            \
0262     CLS& operator=(CLS const&) = delete; \
0263     CLS(CLS&&) = delete;                 \
0264     CLS& operator=(CLS&&) = delete
0265
0266 /*!
0267  * \def CELER_DEFAULT_MOVE_DELETE_COPY
0268  *
0269  * Explicitly declare defaulted copy and move constructors and assignment
0270  * operators. Use this if the destructor is declared explicitly.
0271  */
0272 #define CELER_DEFAULT_MOVE_DELETE_COPY(CLS) \
0273     CLS(CLS const&) = delete;               \
0274     CLS& operator=(CLS const&) = delete;    \
0275     CLS(CLS&&) = default;                   \
0276     CLS& operator=(CLS&&) = default
0277
0278 //!@}
0279 //---------------------------------------------------------------------------//
0280
0281 /*!
0282  * \def CELER_DISCARD
0283  *
0284  * The argument is an unevaluated operand which will generate no code but force
0285  * the expression to be used. This is used in place of the \code
0286  * [[maybe_unused]] \endcode attribute, which actually generates warnings in
0287  * older versions of GCC.
0288  */
0289 #define CELER_DISCARD(CODE) static_cast<void>(sizeof(CODE));
0290
0291 //---------------------------------------------------------------------------//