Merge pull request #183 from LLNL/feature/burmark1/error_macro

Helpers to Print Function Args on Error

Author: MrBurmark

.gitignore

@@ -21,3 +21,8 @@ cmake-build-*/
 docker-*-build/
 /Debug/
 ./extern/googletest/*
+
+# install directories
+install/
+install-*/
+install_*/

docs/sphinx/user_guide/using_camp.rst

@@ -106,6 +106,19 @@ Learn more about the Camp ``tuple`` feature on the :ref:`tuple-label` page.
 
 See the full example `here <https://github.com/LLNL/RAJA/blob/0aef7cc44348d82e94e73e12f77c27ea306e47b8/include/RAJA/util/View.hpp>`_.
 
+RAJA also uses Camp for error checking and generating error strings for cuda
+and hip API functions.
+
+.. code-block:: cpp
+
+   CAMP_HIP_API_INVOKE_AND_CHECK(hipLaunchKernel, func, gridDim, blockDim, args, shmem, res.get_stream());
+   // C++ exception with description "HIP error: invalid configuration argument hipLaunchKernel(func=0x273ff0, gridDim={1,2,3}, blockDim={3,2,1}, args=0x7fffffff76f8, sharedMem=0, stream=0x7e0860) /path/to/RAJA/install/RAJA/include/RAJA/policy/hip/MemUtils_HIP.hpp:273" thrown in the test body.
+
+In this example, RAJA uses a CAMP error checking and reporting macro when
+launching a hip kernel. If the kernel launch fails it will generate a string
+containing the error message, the function called, the function arguments, and
+the location of the call.
+
 RAJA has many examples of using Camp. In fact, so many internal RAJA implementations use Camp that RAJA has a
 `Camp alias page <https://github.com/LLNL/RAJA/blob/0aef7cc44348d82e94e73e12f77c27ea306e47b8/include/RAJA/util/camp_aliases.hpp>`_ which
 creates RAJA aliases for many Camp features.

include/camp/defines.hpp

@@ -10,6 +10,7 @@
 
 #include <cstddef>
 #include <cstdint>
+#include <string>
 
 #include <camp/config.hpp>
 
@@ -29,6 +30,9 @@ namespace camp
 
 #define CAMP_ALLOW_UNUSED_LOCAL(X) (void)(X)
 
+#define CAMP_STRINGIFY_HELPER(...) #__VA_ARGS__
+#define CAMP_STRINGIFY(...) CAMP_STRINGIFY_HELPER(__VA_ARGS__)
+
 #if defined(__clang__)
 #define CAMP_COMPILER_CLANG
 #elif defined(__INTEL_COMPILER)
@@ -175,19 +179,58 @@ using nullptr_t = decltype(nullptr);
     using type = typename X<Lambda::template expr, Rest...>::type; \
   }
 
+
 /// Throw a runtime_error, avoid including exception everywhere
+[[noreturn]]
 CAMP_DLL_EXPORT void throw_re(const char *s);
+///
+[[noreturn]]
+CAMP_DLL_EXPORT void throw_re(std::string const& s);
 
 #ifdef CAMP_ENABLE_CUDA
 
 #define campCudaErrchk(ans) ::camp::cudaAssert((ans), #ans, __FILE__, __LINE__)
 
 #define campCudaErrchkDiscardReturn(ans) (void)::camp::cudaAssert((ans), #ans, __FILE__, __LINE__)
 
+[[deprecated]]
 CAMP_DLL_EXPORT cudaError_t cudaAssert(cudaError_t code,
-                              const char *call,
-                              const char *file,
-                              int line);
+                                       const char *call,
+                                       const char *file,
+                                       int line);
+
+/// Invoke the given CUDA API function and report CUDA error codes.
+///
+/// NOTE: This prints arguments to the given function so the arguments
+///       are evaluated twice on an error.
+#define CAMP_CUDA_API_INVOKE_AND_CHECK_IMPLEMENTATION(func, ...)             \
+    /* Avoid shadowing by adding 56792578 to variable names */               \
+    cudaError_t code_56792578 = func(__VA_ARGS__);                           \
+    if (code_56792578 != cudaSuccess &&                                      \
+        code_56792578 != cudaErrorNotReady) /* [[unlikely]] */               \
+    {                                                                        \
+      static constexpr auto func_name_56792578 = CAMP_STRINGIFY(func);       \
+      static constexpr auto arg_names_56792578 =                             \
+          ::camp::experimental::cuda_get_api_arg_names(func_name_56792578);  \
+      ::camp::reportError(                                                   \
+          "CUDA",                                                            \
+          cudaGetErrorString(code_56792578),                                 \
+          func_name_56792578,                                                \
+          arg_names_56792578,                                                \
+          std::forward_as_tuple(__VA_ARGS__),                                \
+          __FILE__, __LINE__);                                               \
+    }
+///
+#define CAMP_CUDA_API_INVOKE_AND_CHECK_RETURN(...)                           \
+  [&](){                                                                     \
+    CAMP_CUDA_API_INVOKE_AND_CHECK_IMPLEMENTATION(__VA_ARGS__)               \
+    return code_56792578;                                                    \
+  }()
+///
+#define CAMP_CUDA_API_INVOKE_AND_CHECK(...)                                  \
+  do {                                                                       \
+    CAMP_CUDA_API_INVOKE_AND_CHECK_IMPLEMENTATION(__VA_ARGS__)               \
+  } while (false)
 
 #endif  //#ifdef CAMP_ENABLE_CUDA
 
@@ -198,10 +241,45 @@ CAMP_DLL_EXPORT cudaError_t cudaAssert(cudaError_t code,
 
 #define campHipErrchkDiscardReturn(ans) (void)::camp::hipAssert((ans), #ans, __FILE__, __LINE__)
 
+[[deprecated]]
 CAMP_DLL_EXPORT hipError_t hipAssert(hipError_t code,
-                            const char *call,
-                            const char *file,
-                            int line);
+                                     const char *call,
+                                     const char *file,
+                                     int line);
+
+
+/// Invoke the given HIP API function and report HIP error codes.
+///
+/// NOTE: This prints arguments to the given function so the arguments
+///       are evaluated twice on an error.
+#define CAMP_HIP_API_INVOKE_AND_CHECK_IMPLEMENTATION(func, ...)              \
+    /* Avoid shadowing by adding 56792578 to variable names */               \
+    hipError_t code_56792578 = func(__VA_ARGS__);                            \
+    if (code_56792578 != hipSuccess &&                                       \
+        code_56792578 != hipErrorNotReady) /* [[unlikely]] */                \
+    {                                                                        \
+      static constexpr auto func_name_56792578 = CAMP_STRINGIFY(func);       \
+      static constexpr auto arg_names_56792578 =                             \
+          ::camp::experimental::hip_get_api_arg_names(func_name_56792578);   \
+      ::camp::reportError(                                                   \
+          "HIP",                                                             \
+          hipGetErrorString(code_56792578),                                  \
+          func_name_56792578,                                                \
+          arg_names_56792578,                                                \
+          std::forward_as_tuple(__VA_ARGS__),                                \
+          __FILE__, __LINE__);                                               \
+    }
+///
+#define CAMP_HIP_API_INVOKE_AND_CHECK_RETURN(...)                            \
+  [&](){                                                                     \
+    CAMP_HIP_API_INVOKE_AND_CHECK_IMPLEMENTATION(__VA_ARGS__)                \
+    return code_56792578;                                                    \
+  }()
+///
+#define CAMP_HIP_API_INVOKE_AND_CHECK(...)                                   \
+  do {                                                                       \
+    CAMP_HIP_API_INVOKE_AND_CHECK_IMPLEMENTATION(__VA_ARGS__)                \
+  } while (false)
 
 #endif  //#ifdef CAMP_ENABLE_HIP
 

include/camp/helpers.hpp

@@ -10,7 +10,20 @@
 
 #include <cstddef>
 #include <utility>
+#include <cstddef>
+#include <cstdint>
+#include <iostream>
+#include <sstream>
+#include <string>
+#include <string_view>
+#include <array>
+#include <tuple>
+#include <utility>
+#include <algorithm>
+#include <stdexcept>
+#include <type_traits>
 
+#include "camp/config.hpp"
 #include "camp/defines.hpp"
 
 namespace camp
@@ -228,6 +241,211 @@ CAMP_HOST_DEVICE void safe_swap(T& t1, T& t2)
   using std::swap;
   swap(t1, t2);
 }
+
+
+namespace experimental
+{
+
+//! Printing helper class to add printability to types defined outside of camp.
+//
+//  Specialize to customize printing or add printing for a type. This avoids
+//  conflicts if an operator<< is later added for the type in question.
+//  Write specializations for non-const reference and const-reference.
+template < typename T >
+struct StreamInsertHelper
+{
+  static_assert(std::is_lvalue_reference_v<T>);
+
+  T m_val;
+
+  std::ostream& operator()(std::ostream& str) const
+  {
+    return str << m_val;
+  }
+};
+
+// deduction guide for StreamInsertHelper
+template<typename T>
+StreamInsertHelper(T& val) -> StreamInsertHelper<T&>;
+template<typename T>
+StreamInsertHelper(T const& val) -> StreamInsertHelper<T const&>;
+template<typename T>
+StreamInsertHelper(T&& val) -> StreamInsertHelper<T const&>;
+template<typename T>
+StreamInsertHelper(T const&& val) -> StreamInsertHelper<T const&>;
+
+// Allow printing of StreamInsertHelper using its call operator
+template < typename T >
+inline std::ostream& operator<<(std::ostream& str, StreamInsertHelper<T> const& si)
+{
+  return si(str);
+}
+
+
+#ifdef CAMP_ENABLE_CUDA
+
+//! Get the argument names for the given cuda API function name.
+//
+//  Returns a space separated string of the arguments to the given function.
+//  Returns an empty string if func is unknown.
+constexpr std::string_view cuda_get_api_arg_names(std::string_view func)
+{
+  using storage_type = std::pair<const char*, const char*>;
+  constexpr storage_type known_functions[] {
+    {"cudaDeviceSynchronize",                         ""},
+    {"cudaGetDevice",                                 "device"},
+    {"cudaGetDeviceProperties",                       "prop device"},
+    {"cudaSetDevice",                                 "device"},
+    {"cudaStreamCreate",                              "pStream"},
+    {"cudaStreamSynchronize",                         "stream"},
+    {"cudaStreamWaitEvent",                           "stream event flags"},
+    {"cudaEventCreateWithFlags",                      "event flags"},
+    {"cudaEventSynchronize",                          "event"},
+    {"cudaEventQuery",                                "event"},
+    {"cudaEventRecord",                               "event stream"},
+    {"cudaHostAlloc",                                 "pHost size flags"},
+    {"cudaMallocHost",                                "ptr size"},
+    {"cudaMalloc",                                    "devPtr size"},
+    {"cudaMallocManaged",                             "devPtr size flags"},
+    {"cudaHostFree",                                  "ptr"},
+    {"cudaFree",                                      "devPtr"},
+    {"cudaMemset",                                    "devPtr value count"},
+    {"cudaMemcpy",                                    "dst src count kind"},
+    {"cudaMemsetAsync",                               "devPtr value count stream"},
+    {"cudaMemcpyAsync",                               "dst src count kind stream"},
+    {"cudaLaunchKernel",                              "func gridDim blockDim args sharedMem stream"},
+    {"cudaPeekAtLastError",                           ""},
+    {"cudaGetLastError",                              ""},
+    {"cudaFuncGetAttributes",                         "attr func"},
+    {"cudaOccupancyMaxPotentialBlockSize",            "minGridSize blockSize func dynamicSMemSize blockSizeLimit"},
+    {"cudaOccupancyMaxActiveBlocksPerMultiprocessor", "numBlocks func blockSize dynamicSMemSize"}
+  };
+  for (auto [api_name, api_args] : known_functions) {
+    if (func == api_name) {
+      return api_args;
+    }
+  }
+  return "";
+}
+
+#endif  //#ifdef CAMP_ENABLE_CUDA
+
+#ifdef CAMP_ENABLE_HIP
+
+//! Get the argument names for the given hip API function name.
+//
+//  Returns a space separated string of the arguments to the given function.
+//  Returns an empty string if func is unknown.
+constexpr std::string_view hip_get_api_arg_names(std::string_view func)
+{
+  using storage_type = std::pair<const char*, const char*>;
+  constexpr storage_type known_functions[] {
+    {"hipDeviceSynchronize",                         ""},
+    {"hipGetDevice",                                 "device"},
+    {"hipGetDeviceProperties",                       "prop device"},
+    {"hipGetDevicePropertiesR0600",                  "prop device"},
+    {"hipSetDevice",                                 "device"},
+    {"hipStreamCreate",                              "pStream"},
+    {"hipStreamSynchronize",                         "stream"},
+    {"hipStreamWaitEvent",                           "stream event flags"},
+    {"hipEventCreateWithFlags",                      "event flags"},
+    {"hipEventSynchronize",                          "event"},
+    {"hipEventQuery",                                "event"},
+    {"hipEventRecord",                               "event stream"},
+    {"hipHostMalloc",                                "pHost size flags"},
+    {"hipMalloc",                                    "devPtr size"},
+    {"hipMallocManaged",                             "devPtr size flags"},
+    {"hipHostFree",                                  "ptr"},
+    {"hipFree",                                      "devPtr"},
+    {"hipMemset",                                    "devPtr value count"},
+    {"hipMemcpy",                                    "dst src count kind"},
+    {"hipMemsetAsync",                               "devPtr value count stream"},
+    {"hipMemcpyAsync",                               "dst src count kind stream"},
+    {"hipLaunchKernel",                              "func gridDim blockDim args sharedMem stream"},
+    {"hipPeekAtLastError",                           ""},
+    {"hipGetLastError",                              ""},
+    {"hipFuncGetAttributes",                         "attr func"},
+    {"hipOccupancyMaxPotentialBlockSize",            "minGridSize blockSize func dynamicSMemSize blockSizeLimit"},
+    {"hipOccupancyMaxActiveBlocksPerMultiprocessor", "numBlocks func blockSize dynamicSMemSize"}
+  };
+  for (auto [api_name, api_args] : known_functions) {
+    if (func == api_name) {
+      return api_args;
+    }
+  }
+  return "";
+}
+
+#endif  //#ifdef CAMP_ENABLE_HIP
+
+
+//! Generate a string for the given args
+//
+//  This function generates a string for the given argument names and arguments.
+//  Uses StreamInsertHelper to stringify the types in args.
+template < typename Tuple, std::size_t ...Is >
+std::ostream& insertArgsString(std::ostream& str,
+                               std::string_view arg_names,
+                               Tuple&& args_tuple,
+                               std::index_sequence<Is...>)
+{
+  auto insert_arg = [&, first=true](auto&& arg) mutable {
+    if (first) {
+      first = false;
+    } else {
+      str << ", ";
+    }
+    if (!arg_names.empty()) {
+      auto arg_name_size = arg_names.find(' ');
+      str << arg_names.substr(0, arg_name_size) << "=";
+      if (arg_name_size < arg_names.size()) {
+        ++arg_name_size; // skip space
+      }
+      arg_names.remove_prefix(arg_name_size);
+    }
+    str << ::camp::experimental::StreamInsertHelper{arg};
+  };
+  ::camp::sink(insert_arg); // suppress unused warning from nvcc
+  using std::get;
+  (..., insert_arg(get<Is>(std::forward<Tuple>(args_tuple))));
+  return str;
+}
+
+}  // namespace experimental
+
+
+/// Report hip errors by throwing an exception or printing to cerr
+///
+/// This function generates an error message by getting a string for the given
+/// hip error code, function, argument names, arguments, and source location
+/// information. Uses StreamInsertHelper to stringify the types in args.
+///
+/// This function throws an exception if abort is true otherwise prints to cerr.
+template < typename Tuple >
+void reportError(std::string_view error_name,
+                 std::string_view error_description,
+                 std::string_view func_name,
+                 std::string_view arg_names,
+                 Tuple&& args_tuple,
+                 std::string_view file,
+                 int line,
+                 bool abort = true)
+{
+  std::ostringstream str;
+  str << error_name << " error: " << error_description;
+  str << " " << func_name << "(";
+  ::camp::experimental::insertArgsString(
+      str, arg_names,
+      std::forward<Tuple>(args_tuple),
+      std::make_index_sequence<std::tuple_size_v<std::decay_t<Tuple>>>{});
+  str << ") " << file << ":" << line;
+  if (abort) {
+    ::camp::throw_re(str.str());
+  } else {
+    std::cerr << str.str();
+  }
+}
+
 }  // namespace camp
 
 #endif /* CAMP_HELPERS_HPP */

include/camp/resource.hpp

@@ -13,6 +13,8 @@
 #include <mutex>
 #include <type_traits>
 
+#include "camp/config.hpp"
+#include "camp/defines.hpp"
 #include "camp/helpers.hpp"
 #include "camp/resource/event.hpp"
 #include "camp/resource/host.hpp"