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 /*!
0074  * \def CELER_CONSTEXPR_FUNCTION
0075  *
0076  * Decorate a function that works on both host and device, with and without
0077  * NVCC, can be evaluated at compile time, and should be forcibly inlined.
0078  */
0079 #define CELER_CONSTEXPR_FUNCTION constexpr CELER_FORCEINLINE_FUNCTION
0080
0081 /*!
0082  * \def CELER_UNLIKELY(condition)
0083  *
0084  * Mark the result of this condition to be "unlikely".
0085  *
0086  * This asks the compiler to move the section of code to a "cold" part of the
0087  * instructions, improving instruction locality. It should be used primarily
0088  * for error checking conditions.
0089  */
0090 #if defined(__clang__) || defined(__GNUC__)
0091 // GCC and Clang support the same builtin
0092 #    define CELER_UNLIKELY(COND) __builtin_expect(!!(COND), 0)
0093 #else
0094 // No other compilers seem to have a similar builtin
0095 #    define CELER_UNLIKELY(COND) (COND)
0096 #endif
0097
0098 /*!
0099  * \def CELER_UNREACHABLE
0100  *
0101  * Mark a point in code as being impossible to reach in normal execution.
0102  *
0103  * See https://clang.llvm.org/docs/LanguageExtensions.html#builtin-unreachable
0104  * or https://msdn.microsoft.com/en-us/library/1b3fsfxw.aspx
0105  * or
0106  * https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#__builtin_unreachable
0107  *
0108  * (The "unreachable" and "assume" compiler optimizations for CUDA are only
0109  * available in API version 11.3 or higher, which is encoded as
0110  * \code major*1000 + minor*10 \endcode).
0111  *
0112  * \note This macro should not generally be used; instead, the macro \c
0113  * CELER_ASSERT_UNREACHABLE() defined in base/Assert.hh should be used instead
0114  * (to provide a more detailed error message in case the point *is* reached).
0115  */
0116 #if (!defined(__CUDA_ARCH__) && (defined(__clang__) || defined(__GNUC__))) \
0117     || defined(__NVCOMPILER)                                               \
0118     || (defined(__CUDA_ARCH__) && CUDART_VERSION >= 11030)                 \
0119     || defined(__HIP_DEVICE_COMPILE__)
0120 #    define CELER_UNREACHABLE __builtin_unreachable()
0121 #elif defined(_MSC_VER)
0122 #    define CELER_UNREACHABLE __assume(false)
0123 #else
0124 #    define CELER_UNREACHABLE
0125 #endif
0126
0127 /*!
0128  * \def CELER_USE_DEVICE
0129  *
0130  * True if HIP or CUDA are enabled, false otherwise.
0131  */
0132 #if CELERITAS_USE_CUDA || CELERITAS_USE_HIP
0133 #    define CELER_USE_DEVICE 1
0134 #else
0135 #    define CELER_USE_DEVICE 0
0136 #endif
0137
0138 /*!
0139  * \def CELER_DEVICE_SOURCE
0140  *
0141  * Defined and true if building a HIP or CUDA source file. This is a generic
0142  * replacement for \c __CUDACC__ .
0143  */
0144 #if defined(__CUDACC__) || defined(__HIP__)
0145 #    define CELER_DEVICE_SOURCE 1
0146 #elif defined(__DOXYGEN__)
0147 #    define CELER_DEVICE_SOURCE 0
0148 #endif
0149
0150 /*!
0151  * \def CELER_DEVICE_COMPILE
0152  *
0153  * Defined and true if building device code in HIP or CUDA. This is a generic
0154  * replacement for \c __CUDA_ARCH__ .
0155  */
0156 #if defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__)
0157 #    define CELER_DEVICE_COMPILE 1
0158 #elif defined(__DOXYGEN__)
0159 #    define CELER_DEVICE_COMPILE 0
0160 #endif
0161
0162 #if CELERITAS_USE_CUDA            \
0163     && (__CUDACC_VER_MAJOR__ < 11 \
0164         || (__CUDACC_VER_MAJOR__ == 11 && __CUDACC_VER_MINOR__ < 5))
0165 /*!
0166  * Work around older NVCC bugs with `if constexpr`.
0167  *
0168  * These cause errors such as \verbatim
0169  *    error: missing return statement at end of non-void function
0170  * \endverbatim
0171  */
0172 #    define CELER_CUDACC_BUGGY_IF_CONSTEXPR 1
0173 #else
0174 #    define CELER_CUDACC_BUGGY_IF_CONSTEXPR 0
0175 #endif
0176
0177 //!@}
0178 //---------------------------------------------------------------------------//
0179 //!@{
0180 //! \name Exception handling macros
0181
0182 /*!
0183  * \def CELER_TRY_HANDLE
0184  *
0185  * "Try" to execute the statement, and "handle" *all* thrown errors by calling
0186  * the given function-like error handler with a \c std::exception_ptr object.
0187  *
0188  * \note A file that uses this macro must include the \c \<exception\> header
0189  * (but since the \c HANDLE_EXCEPTION needs to take an exception pointer, it's
0190  * got to be included anyway).
0191  */
0192 #define CELER_TRY_HANDLE(STATEMENT, HANDLE_EXCEPTION)   \
0193     do                                                  \
0194     {                                                   \
0195         try                                             \
0196         {                                               \
0197             STATEMENT;                                  \
0198         }                                               \
0199         catch (...)                                     \
0200         {                                               \
0201             HANDLE_EXCEPTION(std::current_exception()); \
0202         }                                               \
0203     } while (0)
0204
0205 /*!
0206  * \def CELER_TRY_HANDLE_CONTEXT
0207  *
0208  * Try the given statement, and if it fails, chain it into the given exception.
0209  *
0210  * The given \c CONTEXT_EXCEPTION must be an expression that yields an rvalue
0211  * to a \c std::exception subclass that isn't \c final . The resulting chained
0212  * exception will be passed into \c HANDLE_EXCEPTION for processing.
0213  */
0214 #define CELER_TRY_HANDLE_CONTEXT(                          \
0215     STATEMENT, HANDLE_EXCEPTION, CONTEXT_EXCEPTION)        \
0216     CELER_TRY_HANDLE(                                      \
0217         do {                                               \
0218             try                                            \
0219             {                                              \
0220                 STATEMENT;                                 \
0221             }                                              \
0222             catch (...)                                    \
0223             {                                              \
0224                 std::throw_with_nested(CONTEXT_EXCEPTION); \
0225             }                                              \
0226         } while (0),                                       \
0227         HANDLE_EXCEPTION)
0228
0229 //!@}
0230 //---------------------------------------------------------------------------//
0231 //!@{
0232 //! \name Class definition macros
0233
0234 /*!
0235  * \def CELER_DEFAULT_COPY_MOVE
0236  *
0237  * Explicitly declare defaulted copy and move constructors and assignment
0238  * operators. Use this if the destructor is declared explicitly, or as part of
0239  * the "protected" section of an interface class to prevent assignment between
0240  * incompatible classes.
0241  */
0242 #define CELER_DEFAULT_COPY_MOVE(CLS)      \
0243     CLS(CLS const&) = default;            \
0244     CLS& operator=(CLS const&) = default; \
0245     CLS(CLS&&) = default;                 \
0246     CLS& operator=(CLS&&) = default
0247
0248 /*!
0249  * \def CELER_DELETE_COPY_MOVE
0250  *
0251  * Explicitly declare *deleted* copy and move constructors and assignment
0252  * operators. Use this for scoped RAII classes.
0253  */
0254 #define CELER_DELETE_COPY_MOVE(CLS)      \
0255     CLS(CLS const&) = delete;            \
0256     CLS& operator=(CLS const&) = delete; \
0257     CLS(CLS&&) = delete;                 \
0258     CLS& operator=(CLS&&) = delete
0259
0260 /*!
0261  * \def CELER_DEFAULT_MOVE_DELETE_COPY
0262  *
0263  * Explicitly declare defaulted copy and move constructors and assignment
0264  * operators. Use this if the destructor is declared explicitly.
0265  */
0266 #define CELER_DEFAULT_MOVE_DELETE_COPY(CLS) \
0267     CLS(CLS const&) = delete;               \
0268     CLS& operator=(CLS const&) = delete;    \
0269     CLS(CLS&&) = default;                   \
0270     CLS& operator=(CLS&&) = default
0271
0272 //!@}
0273 //---------------------------------------------------------------------------//
0274
0275 /*!
0276  * \def CELER_DISCARD
0277  *
0278  * The argument is an unevaluated operand which will generate no code but force
0279  * the expression to be used. This is used in place of the \code
0280  * [[maybe_unused]] \endcode attribute, which actually generates warnings in
0281  * older versions of GCC.
0282  */
0283 #define CELER_DISCARD(CODE) static_cast<void>(sizeof(CODE));
0284
0285 //---------------------------------------------------------------------------//