VecGeom/base/Assert.h

0001 //------------------------------- -*- C++ -*- -------------------------------//
0002 // Copyright Celeritas contributors: see https://github.com/celeritas-project/celeritas?tab=License-1-ov-file
0003 // SPDX-License-Identifier: (Apache-2.0 OR MIT)
0004 //---------------------------------------------------------------------------//
0005 /*!
0006  * \file Assert.h
0007  * \brief Macros, exceptions, and helpers for assertions and error handling.
0008  *
0009  * This defines host- and device-compatible assertion macros that are
0010  * currenntly toggled on the \c NDEBUG configure macros.
0011  *
0012  * Derived from celeritas corecel: DeviceRuntimeApi.hh, Macros.hh, Assert.hh
0013  */
0014 //---------------------------------------------------------------------------/
0015
0016 #ifndef VECGEOM_BASE_ASSERT_H_
0017 #define VECGEOM_BASE_ASSERT_H_
0018
0019 #include "VecGeom/base/Config.h"
0020 #include "VecGeom/base/Global.h"
0021
0022 #ifndef __CUDA_ARCH__
0023 #include <stdexcept>
0024 #include <sstream>
0025 #else
0026 #include <cassert>
0027 #endif
0028
0029 // NOTE: if we want more debug granularity, these could be defined via CMake config in the future
0030 #ifdef NDEBUG
0031 #define VECGEOM_DEBUG 0
0032 #define VECGEOM_DEVICE_DEBUG 0
0033 #else
0034 #define VECGEOM_DEBUG 1
0035 #define VECGEOM_DEVICE_DEBUG 1
0036 #endif
0037
0038 /*!
0039  * \def VECGEOM_DEVICE_PLATFORM
0040  *
0041  * API prefix token for the device offloading type.
0042  */
0043 /*!
0044  * \def VECGEOM_DEVICE_API_SYMBOL
0045  *
0046  * Add a prefix "hip" or "cuda" to a code token.
0047  */
0048 #if defined(VECGEOM_ENABLE_CUDA)
0049 #define VECGEOM_DEVICE_PLATFORM cuda
0050 #define VECGEOM_DEVICE_PLATFORM_UPPER_STR "CUDA"
0051 #define VECGEOM_DEVICE_API_SYMBOL(TOK) cuda##TOK
0052 #elif defined(VECGEOM_ENABLE_HIP)
0053 // NOTE: not yet implemented
0054 #define VECGEOM_DEVICE_PLATFORM hip
0055 #define VECGEOM_DEVICE_PLATFORM_UPPER_STR "HIP"
0056 #define VECGEOM_DEVICE_API_SYMBOL(TOK) hip##TOK
0057 #else
0058 #define VECGEOM_DEVICE_PLATFORM none
0059 #define VECGEOM_DEVICE_PLATFORM_UPPER_STR ""
0060 #define VECGEOM_DEVICE_API_SYMBOL(TOK) void
0061 #endif
0062
0063 /*!
0064  * \def VECGEOM_DEVICE_SOURCE
0065  *
0066  * Defined and true if building a HIP or CUDA source file. This is a generic
0067  * replacement for \c __CUDACC__ .
0068  */
0069 /*!
0070  * \def VECGEOM_DEVICE_COMPILE
0071  *
0072  * Defined and true if building device code in HIP or CUDA. This is a generic
0073  * replacement for \c __CUDA_ARCH__ .
0074  */
0075 #if defined(__CUDACC__) || defined(__HIP__)
0076 #define VECGEOM_DEVICE_SOURCE 1
0077 #elif defined(__DOXYGEN__)
0078 #define VECGEOM_DEVICE_SOURCE 0
0079 #endif
0080
0081 #if defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__)
0082 #define VECGEOM_DEVICE_COMPILE 1
0083 #elif defined(__DOXYGEN__)
0084 #define VECGEOM_DEVICE_COMPILE 0
0085 #endif
0086
0087 /*!
0088  * \def VECGEOM_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 VECGEOM_UNLIKELY(COND) __builtin_expect(!!(COND), 0)
0099 #else
0100 // No other compilers seem to have a similar builtin
0101 #define VECGEOM_UNLIKELY(COND) (COND)
0102 #endif
0103
0104 /*!
0105  * \def VECGEOM_ASSERT
0106  *
0107  * Internal debug assertion macro. This replaces standard \c assert usage.
0108  */
0109 /*!
0110  * \def VECGEOM_NOT_CONFIGURED
0111  *
0112  * Assert if the code point is reached because an optional feature is disabled.
0113  * This generally should be used for the constructors of dummy class
0114  * definitions in, e.g., \c Foo.nocuda.cc:
0115  * \code
0116     Foo::Foo()
0117     {
0118         VECGEOM_NOT_CONFIGURED("CUDA");
0119     }
0120  * \endcode
0121  */
0122 /*!
0123  * \def VECGEOM_DEBUG_FAIL
0124  *
0125  * Throw a debug assertion regardless of the \c VECGEOM_DEBUG setting. This
0126  * is used internally but is also useful for catching subtle programming errors
0127  * in downstream code.
0128  */
0129 /*!
0130  * \def VECGEOM_ASSERT_UNREACHABLE
0131  *
0132  * Throw an assertion if the code point is reached. When debug assertions are
0133  * turned off, this changes to a compiler hint that improves optimization (and
0134  * may force the code to exit uncermoniously if the point is encountered,
0135  * rather than continuing on with undefined behavior).
0136  */
0137 /*!
0138  * \def VECGEOM_VALIDATE
0139  *
0140  * Always-on runtime assertion macro. This can check user input and input data
0141  * consistency, and will raise std::runtime_error on failure with a descriptive error
0142  * message that is streamed as the second argument. If used
0143  * in \c __device__ -annotated code, the second argument *must* be a single C string.
0144  *
0145  * An always-on debug-type assertion without a detailed message can be
0146  * constructed by omitting the stream (but leaving the comma):
0147  * \code
0148     VECGEOM_VALIDATE(file_stream,);
0149  * \endcode
0150  */
0151
0152 #if !defined(__HIP__) && !defined(__CUDA_ARCH__)
0153 // Throw in host code
0154 #define VECGEOM_DEBUG_THROW_(MSG, WHICH) throw ::vecgeom::make_debug_error(#WHICH, MSG, __FILE__, __LINE__)
0155 #elif defined(__CUDA_ARCH__) && !defined(NDEBUG)
0156 // Use the assert macro for CUDA when supported
0157 #define VECGEOM_DEBUG_THROW_(MSG, WHICH) assert(false && sizeof(#WHICH ": " MSG))
0158 #else
0159 // Use a special device function to emulate assertion failure if HIP
0160 // (assertion from multiple threads simultaeously can cause unexpected device
0161 // failures on AMD hardware) or if NDEBUG is in use with CUDA
0162 #define VECGEOM_DEBUG_THROW_(MSG, WHICH) ::vecgeom::device_debug_error(#WHICH, MSG, __FILE__, __LINE__)
0163 #endif
0164
0165 #define VECGEOM_DEBUG_FAIL(MSG, WHICH) \
0166   do {                                 \
0167     VECGEOM_DEBUG_THROW_(MSG, WHICH);  \
0168     ::vecgeom::unreachable();          \
0169   } while (0)
0170 #define VECGEOM_DEBUG_ASSERT_(COND, WHICH) \
0171   do {                                     \
0172     if (VECGEOM_UNLIKELY(!(COND))) {       \
0173       VECGEOM_DEBUG_THROW_(#COND, WHICH);  \
0174     }                                      \
0175   } while (0)
0176 #define VECGEOM_NOASSERT_(COND) \
0177   do {                          \
0178     if (false && (COND)) {      \
0179     }                           \
0180   } while (0)
0181
0182 #if (VECGEOM_DEBUG && !VECGEOM_DEVICE_COMPILE) || (VECGEOM_DEVICE_DEBUG && VECGEOM_DEVICE_COMPILE)
0183 #define VECGEOM_ASSERT(COND) VECGEOM_DEBUG_ASSERT_(COND, internal)
0184 #define VECGEOM_ASSERT_UNREACHABLE() VECGEOM_DEBUG_FAIL("unreachable code point encountered", unreachable)
0185 #else
0186 #define VECGEOM_ASSERT(COND) VECGEOM_NOASSERT_(COND)
0187 #define VECGEOM_ASSERT_UNREACHABLE() ::vecgeom::unreachable()
0188 #endif
0189
0190 #if !VECGEOM_DEVICE_COMPILE
0191 #define VECGEOM_RUNTIME_THROW(WHICH, WHAT, COND) \
0192   throw ::vecgeom::make_runtime_error(WHICH, WHAT, COND, __FILE__, __LINE__)
0193 #else
0194 #define VECGEOM_RUNTIME_THROW(WHICH, WHAT, COND) \
0195   VECGEOM_DEBUG_FAIL("Runtime errors cannot be thrown from device code", unreachable);
0196 #endif
0197
0198 #if !VECGEOM_DEVICE_COMPILE
0199 #define VECGEOM_VALIDATE(COND, MSG)                                           \
0200   do {                                                                        \
0201     if (VECGEOM_UNLIKELY(!(COND))) {                                          \
0202       std::ostringstream vg_runtime_msg_;                                     \
0203       vg_runtime_msg_ MSG;                                                    \
0204       VECGEOM_RUNTIME_THROW("runtime", vg_runtime_msg_.str().c_str(), #COND); \
0205     }                                                                         \
0206   } while (0)
0207 #else
0208 #define VECGEOM_VALIDATE(COND, MSG)                                                            \
0209   do {                                                                                         \
0210     if (VECGEOM_UNLIKELY(!(COND))) {                                                           \
0211       VECGEOM_RUNTIME_THROW("runtime", (::vecgeom::detail::StreamlikeIdentity {} MSG), #COND); \
0212     }                                                                                          \
0213   } while (0)
0214 #endif
0215
0216 #define VECGEOM_NOT_CONFIGURED(WHAT) VECGEOM_RUNTIME_THROW("not configured", WHAT, nullptr)
0217
0218 /*!
0219  * \def VECGEOM_DEVICE_API_CALL
0220  *
0221  * Safely and portably dispatch a CUDA/HIP API call.
0222  *
0223  * When CUDA or HIP support is enabled, execute the wrapped statement
0224  * prepend the argument with "cuda" or "hip" and throw a
0225  * std::runtime_error if it fails. If no device platform is enabled, throw an
0226  * unconfigured assertion.
0227  *
0228  * Example:
0229  *
0230  * \code
0231    VECGEOM_DEVICE_API_CALL(Malloc(&ptr_gpu, 100 * sizeof(float)));
0232    VECGEOM_DEVICE_API_CALL(DeviceSynchronize());
0233  * \endcode
0234  */
0235 #if defined(VECGEOM_ENABLE_CUDA) || defined(VECGEOM_ENABLE_HIP)
0236 #define VECGEOM_DEVICE_API_CALL(STMT)                                                                              \
0237   do {                                                                                                             \
0238     using ErrT_   = VECGEOM_DEVICE_API_SYMBOL(Error_t);                                                            \
0239     ErrT_ result_ = VECGEOM_DEVICE_API_SYMBOL(STMT);                                                               \
0240     if (VECGEOM_UNLIKELY(result_ != VECGEOM_DEVICE_API_SYMBOL(Success))) {                                         \
0241       result_ = VECGEOM_DEVICE_API_SYMBOL(GetLastError)();                                                         \
0242       VECGEOM_RUNTIME_THROW(VECGEOM_DEVICE_PLATFORM_UPPER_STR, VECGEOM_DEVICE_API_SYMBOL(GetErrorString)(result_), \
0243                             #STMT);                                                                                \
0244     }                                                                                                              \
0245   } while (0)
0246 #else
0247 #define VECGEOM_DEVICE_API_CALL(STMT)      \
0248   do {                                     \
0249     VECGEOM_NOT_CONFIGURED("CUDA or HIP"); \
0250   } while (0)
0251 #endif
0252
0253 namespace vecgeom {
0254 //---------------------------------------------------------------------------//
0255 // FUNCTION DECLARATIONS
0256 //---------------------------------------------------------------------------//
0257
0258 #ifndef __CUDA_ARCH__
0259 [[nodiscard]] std::logic_error make_debug_error(char const *which, char const *condition, char const *file, int line);
0260
0261 [[nodiscard]] std::runtime_error make_runtime_error(char const *which, char const *what, char const *condition,
0262                                                     char const *file, int line);
0263 #endif
0264
0265 //---------------------------------------------------------------------------//
0266 // INLINE FUNCTION DEFINITIONS
0267 //---------------------------------------------------------------------------//
0268
0269 //! Invoke undefined behavior
0270 [[noreturn]] inline VECCORE_ATT_HOST_DEVICE void unreachable()
0271 {
0272 #if (!defined(__CUDA_ARCH__) && (defined(__clang__) || defined(__GNUC__))) || defined(__NVCOMPILER) || \
0273     (defined(__CUDA_ARCH__) && CUDART_VERSION >= 11030) || defined(__HIP_DEVICE_COMPILE__)
0274   __builtin_unreachable();
0275 #elif defined(_MSC_VER)
0276   __assume(false);
0277 #else
0278   VECGEOM_UNREACHABLE;
0279 #endif
0280 }
0281
0282 #if defined(__CUDA_ARCH__) && defined(NDEBUG)
0283 //! Host+device definition for CUDA when \c assert is unavailable
0284 inline __attribute__((noinline)) __host__ __device__ void device_debug_error(char const *, char const *condition,
0285                                                                              char const *file, int line)
0286 {
0287   printf("%s:%u:\nvecgeom: internal assertion failed: %s\n", file, line, condition);
0288   __trap();
0289 }
0290 #elif defined(__HIP__)
0291 //! Host-only HIP call (whether or not NDEBUG is in use)
0292 inline __host__ void device_debug_error(char const *which, char const *condition, char const *file, int line)
0293 {
0294   throw make_debug_error(which, condition, file, line);
0295 }
0296
0297 //! Device-only call for HIP (must always be declared; only used if
0298 //! NDEBUG)
0299 inline __attribute__((noinline)) __device__ void device_debug_error(char const *, char const *condition,
0300                                                                     char const *file, int line)
0301 {
0302   printf("%s:%u:\nvecgeom: internal assertion failed: %s\n", file, line, condition);
0303   abort();
0304 }
0305 #endif
0306
0307 namespace detail {
0308 //! Allow passing a single string into a streamlike operator for device-compatible VECGEOM_VALIDATE messages
0309 struct StreamlikeIdentity {
0310    VECCORE_ATT_HOST_DEVICE operator char const *() const { return ""; }
0311 };
0312 inline VECCORE_ATT_HOST_DEVICE char const *operator<<(StreamlikeIdentity const &, char const *s) { return s; }
0313 } // namespace detail
0314
0315 } // namespace vecgeom
0316
0317 #endif