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