codekingpro/portable-devtools

1/******************************************************************************2 * Copyright (c) 2011, Duane Merrill.  All rights reserved.3 * Copyright (c) 2011-2022, NVIDIA CORPORATION.  All rights reserved.4 *5 * Redistribution and use in source and binary forms, with or without6 * modification, are permitted provided that the following conditions are met:7 *     * Redistributions of source code must retain the above copyright8 *       notice, this list of conditions and the following disclaimer.9 *     * Redistributions in binary form must reproduce the above copyright10 *       notice, this list of conditions and the following disclaimer in the11 *       documentation and/or other materials provided with the distribution.12 *     * Neither the name of the NVIDIA CORPORATION nor the13 *       names of its contributors may be used to endorse or promote products14 *       derived from this software without specific prior written permission.15 *16 * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND17 * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED18 * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE19 * DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY20 * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES21 * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;22 * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND23 * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT24 * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS25 * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.26 *27 ******************************************************************************/28 29/**30 * \file31 * Error and event logging routines.32 *33 * The following macros definitions are supported:34 * - \p CUB_LOG.  Simple event messages are printed to \p stdout.35 */36 37#pragma once38 39#include <cub/config.cuh>40 41#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC)42#  pragma GCC system_header43#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG)44#  pragma clang system_header45#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC)46#  pragma system_header47#endif // no system header48 49#include <nv/target>50 51#include <cstdio>52 53CUB_NAMESPACE_BEGIN54 55 56#ifdef DOXYGEN_SHOULD_SKIP_THIS // Only parse this during doxygen passes:57 58/**59 * @def CUB_DEBUG_LOG60 *61 * Causes kernel launch configurations to be printed to the console62 */63#define CUB_DEBUG_LOG64 65/**66 * @def CUB_DEBUG_SYNC67 *68 * Causes synchronization of the stream after every kernel launch to check69 * for errors. Also causes kernel launch configurations to be printed to the70 * console.71 */72#define CUB_DEBUG_SYNC73 74/**75 * @def CUB_DEBUG_HOST_ASSERTIONS76 *77 * Extends `CUB_DEBUG_SYNC` effects by checking host-side precondition78 * assertions.79 */80#define CUB_DEBUG_HOST_ASSERTIONS81 82/**83 * @def CUB_DEBUG_DEVICE_ASSERTIONS84 *85 * Extends `CUB_DEBUG_HOST_ASSERTIONS` effects by checking device-side86 * precondition assertions.87 */88#define CUB_DEBUG_DEVICE_ASSERTIONS89 90/**91 * @def CUB_DEBUG_ALL92 *93 * Causes host and device-side precondition assertions to be checked. Apart94 * from that, causes synchronization of the stream after every kernel launch to95 * check for errors. Also causes kernel launch configurations to be printed to96 * the console.97 */98#define CUB_DEBUG_ALL99 100#endif // DOXYGEN_SHOULD_SKIP_THIS101 102/**103 * \addtogroup UtilMgmt104 * @{105 */106 107 108// `CUB_DETAIL_DEBUG_LEVEL_*`: Implementation details, internal use only:109 110#define CUB_DETAIL_DEBUG_LEVEL_NONE 0111#define CUB_DETAIL_DEBUG_LEVEL_HOST_ASSERTIONS_ONLY 1112#define CUB_DETAIL_DEBUG_LEVEL_LOG 2113#define CUB_DETAIL_DEBUG_LEVEL_SYNC 3114#define CUB_DETAIL_DEBUG_LEVEL_HOST_ASSERTIONS 4115#define CUB_DETAIL_DEBUG_LEVEL_DEVICE_ASSERTIONS 5116#define CUB_DETAIL_DEBUG_LEVEL_ALL 1000117 118// `CUB_DEBUG_*`: User interfaces:119 120// Extra logging, no syncs121#ifdef CUB_DEBUG_LOG122#define CUB_DETAIL_DEBUG_LEVEL CUB_DETAIL_DEBUG_LEVEL_LOG123#endif124 125// Logging + syncs126#ifdef CUB_DEBUG_SYNC127#define CUB_DETAIL_DEBUG_LEVEL CUB_DETAIL_DEBUG_LEVEL_SYNC128#endif129 130// Logging + syncs + host assertions131#ifdef CUB_DEBUG_HOST_ASSERTIONS132#define CUB_DETAIL_DEBUG_LEVEL CUB_DETAIL_DEBUG_LEVEL_HOST_ASSERTIONS133#endif134 135// Logging + syncs + host assertions + device assertions136#ifdef CUB_DEBUG_DEVICE_ASSERTIONS137#define CUB_DETAIL_DEBUG_LEVEL CUB_DETAIL_DEBUG_LEVEL_DEVICE_ASSERTIONS138#endif139 140// All141#ifdef CUB_DEBUG_ALL142#define CUB_DETAIL_DEBUG_LEVEL CUB_DETAIL_DEBUG_LEVEL_ALL143#endif144 145// Default case, no extra debugging:146#ifndef CUB_DETAIL_DEBUG_LEVEL147#ifdef NDEBUG148#define CUB_DETAIL_DEBUG_LEVEL CUB_DETAIL_DEBUG_LEVEL_NONE149#else150#define CUB_DETAIL_DEBUG_LEVEL CUB_DETAIL_DEBUG_LEVEL_HOST_ASSERTIONS_ONLY151#endif152#endif153 154/*155 * `CUB_DETAIL_DEBUG_ENABLE_*`:156 * Internal implementation details, used for testing enabled debug features:157 */158 159#if CUB_DETAIL_DEBUG_LEVEL >= CUB_DETAIL_DEBUG_LEVEL_LOG160#define CUB_DETAIL_DEBUG_ENABLE_LOG161#endif162 163#if CUB_DETAIL_DEBUG_LEVEL >= CUB_DETAIL_DEBUG_LEVEL_SYNC164#define CUB_DETAIL_DEBUG_ENABLE_SYNC165#endif166 167#if (CUB_DETAIL_DEBUG_LEVEL >= CUB_DETAIL_DEBUG_LEVEL_HOST_ASSERTIONS) || \168    (CUB_DETAIL_DEBUG_LEVEL == CUB_DETAIL_DEBUG_LEVEL_HOST_ASSERTIONS_ONLY)169#define CUB_DETAIL_DEBUG_ENABLE_HOST_ASSERTIONS170#endif171 172#if CUB_DETAIL_DEBUG_LEVEL >= CUB_DETAIL_DEBUG_LEVEL_DEVICE_ASSERTIONS173#define CUB_DETAIL_DEBUG_ENABLE_DEVICE_ASSERTIONS174#endif175 176 177/// CUB error reporting macro (prints error messages to stderr)178#if (defined(DEBUG) || defined(_DEBUG)) && !defined(CUB_STDERR)179    #define CUB_STDERR180#endif181 182/**183 * \brief %If \p CUB_STDERR is defined and \p error is not \p cudaSuccess, the184 * corresponding error message is printed to \p stderr (or \p stdout in device185 * code) along with the supplied source context.186 *187 * \return The CUDA error.188 */189__host__ __device__190__forceinline__191cudaError_t Debug(cudaError_t error, const char *filename, int line)192{193  // Clear the global CUDA error state which may have been set by the last194  // call. Otherwise, errors may "leak" to unrelated kernel launches.195 196  // clang-format off197  #ifndef CUB_RDC_ENABLED198  #define CUB_TEMP_DEVICE_CODE199  #else200  #define CUB_TEMP_DEVICE_CODE last_error = cudaGetLastError()201  #endif202 203  cudaError_t last_error = cudaSuccess;204 205  NV_IF_TARGET(206    NV_IS_HOST,207    (last_error = cudaGetLastError();),208    (CUB_TEMP_DEVICE_CODE;)209  );210 211  #undef CUB_TEMP_DEVICE_CODE212  // clang-format on213 214  if (error == cudaSuccess && last_error != cudaSuccess)215  {216    error = last_error;217  }218 219#ifdef CUB_STDERR220  if (error)221  {222    NV_IF_TARGET(223      NV_IS_HOST, (224        fprintf(stderr,225                "CUDA error %d [%s, %d]: %s\n",226                error,227                filename,228                line,229                cudaGetErrorString(error));230        fflush(stderr);231      ),232      (233        printf("CUDA error %d [block (%d,%d,%d) thread (%d,%d,%d), %s, %d]\n",234               error,235               blockIdx.z,236               blockIdx.y,237               blockIdx.x,238               threadIdx.z,239               threadIdx.y,240               threadIdx.x,241               filename,242               line);243      )244    );245  }246#else247  (void)filename;248  (void)line;249#endif250 251  return error;252}253 254/**255 * \brief Debug macro256 */257#ifndef CubDebug258    #define CubDebug(e) CUB_NS_QUALIFIER::Debug((cudaError_t) (e), __FILE__, __LINE__)259#endif260 261 262/**263 * \brief Debug macro with exit264 */265#ifndef CubDebugExit266    #define CubDebugExit(e) if (CUB_NS_QUALIFIER::Debug((cudaError_t) (e), __FILE__, __LINE__)) { exit(1); }267#endif268 269 270/**271 * \brief Log macro for printf statements.272 */273#if !defined(_CubLog)274#if defined(_NVHPC_CUDA) || !(defined(__clang__) && defined(__CUDA__))275 276// NVCC / NVC++277#define _CubLog(format, ...)                                                   \278  do                                                                           \279  {                                                                            \280    NV_IF_TARGET(NV_IS_HOST,                                                   \281                 (printf(format, __VA_ARGS__);),                               \282                 (printf("[block (%d,%d,%d), thread (%d,%d,%d)]: " format,     \283                         blockIdx.z,                                           \284                         blockIdx.y,                                           \285                         blockIdx.x,                                           \286                         threadIdx.z,                                          \287                         threadIdx.y,                                          \288                         threadIdx.x,                                          \289                         __VA_ARGS__);));                                      \290  } while (false)291 292#else // Clang:293 294// XXX shameless hack for clang around variadic printf...295//     Compilies w/o supplying -std=c++11 but shows warning,296//     so we silence them :)297#pragma clang diagnostic ignored "-Wc++11-extensions"298#pragma clang diagnostic ignored "-Wunnamed-type-template-args"299template <class... Args>300inline __host__ __device__ void va_printf(char const *format,301                                          Args const &...args)302{303#ifdef __CUDA_ARCH__304  printf(format,305         blockIdx.z,306         blockIdx.y,307         blockIdx.x,308         threadIdx.z,309         threadIdx.y,310         threadIdx.x,311         args...);312#else313  printf(format, args...);314#endif315}316#ifndef __CUDA_ARCH__317#define _CubLog(format, ...) CUB_NS_QUALIFIER::va_printf(format, __VA_ARGS__);318#else319#define _CubLog(format, ...)                                                   \320  CUB_NS_QUALIFIER::va_printf("[block (%d,%d,%d), thread "                     \321                              "(%d,%d,%d)]: " format,                          \322                              __VA_ARGS__);323#endif324#endif325#endif326 327/** @} */       // end group UtilMgmt328 329CUB_NAMESPACE_END330