GPU: Provide static versions of GPUChkErr() macros

test2

GPU: Provide static versions of GPUChkErr() macros

Author: davidrohr

GPU/Common/CMakeLists.txt

@@ -15,7 +15,7 @@ set(HDRS_INSTALL
     GPUCommonAlgorithm.h
     GPUCommonDef.h
     GPUCommonDefAPI.h
-    GPUCommonChkErr.h
+    GPUCommonHelpers.h
     GPUCommonDefSettings.h
     GPUCommonConstants.h
     GPUCommonLogger.h

GPU/Common/GPUCommonChkErr.h

@@ -35,13 +35,25 @@
     #define GPUCA_GPUCODE // Compiled by GPU compiler
   #endif
 
-  #if defined(__CUDA_ARCH__) || defined(__OPENCL__) || defined(__HIP_DEVICE_COMPILE__)
-    #define GPUCA_GPUCODE_DEVICE // Executed on device
+  #if defined(GPUCA_GPUCODE)
+    #if defined(__CUDA_ARCH__) || defined(__OPENCL__) || defined(__HIP_DEVICE_COMPILE__)
+      #define GPUCA_GPUCODE_DEVICE // Executed on device
+    #endif
+    #if defined(__CUDACC__)
+      #define GPUCA_GPUTYPE CUDA
+    #elif defined(__HIPCC__)
+      #define GPUCA_GPUTYPE HIP
+    #elif defined(__OPENCL__) || defined(__OPENCL_HOST__)
+      #define GPUCA_GPUTYPE OCL
+    #endif
   #endif
 #endif
+#ifndef GPUCA_GPUTYPE
+  #define GPUCA_GPUTYPE CPU
+#endif
 
 #if defined(GPUCA_STANDALONE) || (defined(GPUCA_O2_LIB) && !defined(GPUCA_O2_INTERFACE)) || defined (GPUCA_GPUCODE)
-  #define GPUCA_ALIGPUCODE
+  #define GPUCA_ALIGPUCODE // Part of GPUTracking library but not of interface
 #endif
 
 #if (defined(__CUDACC__) && defined(GPUCA_CUDA_NO_CONSTANT_MEMORY)) || (defined(__HIPCC__) && defined(GPUCA_HIP_NO_CONSTANT_MEMORY)) || (defined(__OPENCL__) && defined(GPUCA_OPENCL_NO_CONSTANT_MEMORY))
@@ -51,13 +63,13 @@
 #endif
 
 #if !defined(GPUCA_GPUCODE) && !defined(GPUCA_STANDALONE) && defined(DEBUG_STREAMER)
-#define GPUCA_DEBUG_STREAMER_CHECK(...) __VA_ARGS__
+  #define GPUCA_DEBUG_STREAMER_CHECK(...) __VA_ARGS__
 #else
-#define GPUCA_DEBUG_STREAMER_CHECK(...)
+  #define GPUCA_DEBUG_STREAMER_CHECK(...)
 #endif
 
 #ifndef GPUCA_RTC_SPECIAL_CODE
-#define GPUCA_RTC_SPECIAL_CODE(...)
+  #define GPUCA_RTC_SPECIAL_CODE(...)
 #endif
 
 // API Definitions for GPU Compilation

GPU/Common/GPUCommonDef.h

@@ -0,0 +1,62 @@
+// Copyright 2019-2020 CERN and copyright holders of ALICE O2.
+// See https://alice-o2.web.cern.ch/copyright for details of the copyright holders.
+// All rights not expressly granted are reserved.
+//
+// This software is distributed under the terms of the GNU General Public
+// License v3 (GPL Version 3), copied verbatim in the file "COPYING".
+//
+// In applying this license CERN does not waive the privileges and immunities
+// granted to it by virtue of its status as an Intergovernmental Organization
+// or submit itself to any jurisdiction.
+
+/// \file GPUCommonHelpers.h
+/// \author David Rohr
+
+// GPUChkErr and GPUChkErrI will both check x for an error, using the loaded backend of GPUReconstruction (requiring GPUReconstruction.h to be included by the user).
+// In case of an error, it will print out the corresponding CUDA / HIP / OpenCL error code
+// GPUChkErr will download GPUReconstruction error values from GPU, print them, and terminate the application with an exception if an error occured.
+// GPUChkErrI will return 0 or 1, depending on whether an error has occurred.
+// These Macros must be called ona GPUReconstruction instance.
+// The GPUChkErrS and GPUChkErrSI are similar but static, without required GPUReconstruction instance.
+// Examples:
+// if (mRec->GPUChkErrI(cudaMalloc(...))) { exit(1); }
+// gpuRecObj.GPUChkErr(cudaMalloc(...));
+// if (GPUChkErrSI(cudaMalloc(..))) { exit(1); }
+
+#ifndef GPUCOMMONHELPERS_H
+#define GPUCOMMONHELPERS_H
+
+// Please #include "GPUReconstruction.h" in your code, if you use these 2!
+#define GPUChkErr(x) GPUChkErrA(x, __FILE__, __LINE__, true)
+#define GPUChkErrI(x) GPUChkErrA(x, __FILE__, __LINE__, false)
+#define GPUChkErrS(x) o2::gpu::internal::GPUReconstructionChkErr(x, __FILE__, __LINE__, true)
+#define GPUChkErrSI(x) o2::gpu::internal::GPUReconstructionChkErr(x, __FILE__, __LINE__, false)
+
+#include "GPUCommonDef.h"
+#include <cstdint>
+
+namespace o2::gpu::internal
+{
+#define GPUCOMMON_INTERNAL_CAT_A(a, b, c) a##b##c
+#define GPUCOMMON_INTERNAL_CAT(...) GPUCOMMON_INTERNAL_CAT_A(__VA_ARGS__)
+extern int32_t GPUCOMMON_INTERNAL_CAT(GPUReconstruction, GPUCA_GPUTYPE, ChkErr)(const int64_t error, const char* file, int32_t line);
+inline int32_t GPUReconstructionCPUChkErr(const int64_t error, const char* file, int32_t line)
+{
+  if (error) {
+    GPUError("GPUCommon Error Code %d (%s:%d)", error, file, line);
+  }
+  return error != 0;
+}
+static inline int32_t GPUReconstructionChkErr(const int64_t error, const char* file, int32_t line, bool failOnError)
+{
+  int32_t retVal = error && GPUCOMMON_INTERNAL_CAT(GPUReconstruction, GPUCA_GPUTYPE, ChkErr)(error, file, line);
+  if (retVal && failOnError) {
+    throw std::runtime_error("GPU API Call Failure");
+  }
+  return error;
+}
+#undef GPUCOMMON_INTERNAL_CAT_A
+#undef GPUCOMMON_INTERNAL_CAT
+} // namespace o2::gpu::internal
+
+#endif

GPU/Common/GPUCommonHelpers.h

@@ -18,7 +18,7 @@ endif()
 message(STATUS "Building GPUTracking with CUDA support ${TMP_TARGET}")
 
 set(SRCS GPUReconstructionCUDA.cu GPUReconstructionCUDAGenRTC.cxx GPUReconstructionCUDAKernels.cu)
-set(HDRS GPUReconstructionCUDA.h GPUReconstructionCUDAInternals.h GPUReconstructionCUDADef.h GPUReconstructionCUDAIncludesHost.h CUDAThrustHelpers.h)
+set(HDRS GPUReconstructionCUDA.h GPUReconstructionCUDAInternals.h GPUReconstructionCUDAHelpers.inc GPUReconstructionCUDADef.h GPUReconstructionCUDAIncludesHost.h CUDAThrustHelpers.h)
 # -------------------------------- Prepare RTC -------------------------------------------------------
 enable_language(ASM)
 if(ALIGPU_BUILD_TYPE STREQUAL "O2")

GPU/GPUTracking/Base/cuda/CMakeLists.txt

@@ -22,6 +22,7 @@
 #include "CUDAThrustHelpers.h"
 #include "GPUReconstructionIncludes.h"
 #include "GPUParamRTC.h"
+#include "GPUReconstructionCUDAHelpers.inc"
 
 #if defined(GPUCA_KERNEL_COMPILE_MODE) && GPUCA_KERNEL_COMPILE_MODE == 1
 #include "utils/qGetLdBinarySymbols.h"
@@ -62,13 +63,9 @@ GPUReconstructionCUDABackend::~GPUReconstructionCUDABackend()
 }
 
 static_assert(sizeof(cudaError_t) <= sizeof(int64_t) && cudaSuccess == 0);
-int32_t GPUReconstructionCUDABackend::GPUChkErrStatic(const int64_t error, const char* file, int32_t line)
+int32_t GPUReconstructionCUDABackend::GPUChkErrInternal(const int64_t error, const char* file, int32_t line) const
 {
-  if (error == cudaSuccess) {
-    return (0);
-  }
-  GPUError("CUDA Error: %ld / %s (%s:%d)", error, cudaGetErrorString((cudaError_t)error), file, line);
-  return 1;
+  return internal::GPUReconstructionCUDAChkErr(error, file, line);
 }
 
 GPUReconstructionCUDA::GPUReconstructionCUDA(const GPUSettingsDeviceBackend& cfg) : GPUReconstructionKernels(cfg)

GPU/GPUTracking/Base/cuda/GPUReconstructionCUDA.cu

@@ -33,13 +33,12 @@ class GPUReconstructionCUDABackend : public GPUReconstructionDeviceBase
 {
  public:
   ~GPUReconstructionCUDABackend() override;
-  static int32_t GPUChkErrStatic(const int64_t error, const char* file, int32_t line);
 
  protected:
   GPUReconstructionCUDABackend(const GPUSettingsDeviceBackend& cfg);
 
   void PrintKernelOccupancies() override;
-  virtual int32_t GPUChkErrInternal(const int64_t error, const char* file, int32_t line) const override { return GPUChkErrStatic(error, file, line); }
+  virtual int32_t GPUChkErrInternal(const int64_t error, const char* file, int32_t line) const override;
 
   template <class T, int32_t I = 0, typename... Args>
   void runKernelBackend(const krnlSetupArgs<T, I, Args...>& args);

GPU/GPUTracking/Base/cuda/GPUReconstructionCUDA.h

@@ -32,11 +32,12 @@ using namespace o2::gpu;
 #include "TrackParametrizationWithError.cxx"
 #include "Propagator.cxx"
 #include "TrackLTIntegral.cxx"
+#include "GPUReconstructionCUDAHelpers.inc"
 
 #ifndef GPUCA_NO_CONSTANT_MEMORY
 static GPUReconstructionDeviceBase::deviceConstantMemRegistration registerConstSymbol([]() {
   void* retVal = nullptr;
-  if (GPUReconstructionCUDA::GPUChkErrStatic(cudaGetSymbolAddress(&retVal, gGPUConstantMemBuffer), __FILE__, __LINE__)) {
+  if (GPUChkErrS(cudaGetSymbolAddress(&retVal, gGPUConstantMemBuffer))) {
     throw std::runtime_error("Could not obtain GPU constant memory symbol");
   }
   return retVal;

GPU/GPUTracking/Base/cuda/GPUReconstructionCUDAExternalProvider.cu

@@ -0,0 +1,31 @@
+// Copyright 2019-2020 CERN and copyright holders of ALICE O2.
+// See https://alice-o2.web.cern.ch/copyright for details of the copyright holders.
+// All rights not expressly granted are reserved.
+//
+// This software is distributed under the terms of the GNU General Public
+// License v3 (GPL Version 3), copied verbatim in the file "COPYING".
+//
+// In applying this license CERN does not waive the privileges and immunities
+// granted to it by virtue of its status as an Intergovernmental Organization
+// or submit itself to any jurisdiction.
+
+/// \file GPUReconstructionCUDAHelpers.inc
+/// \author David Rohr
+
+#ifndef GPURECONSTRUCTIONCUDAHELPERS_INC_H
+#define GPURECONSTRUCTIONCUDAHELPERS_INC_H
+
+#include "GPUCommonHelpers.h"
+
+namespace o2::gpu::internal
+{
+int32_t __attribute__((weak)) GPUReconstructionCUDAChkErr(const int64_t error, const char* file, int32_t line)
+{
+  if (error != cudaSuccess) {
+    GPUError("CUDA Error: %ld / %s (%s:%d)", error, cudaGetErrorString((cudaError_t)error), file, line);
+  }
+  return error != cudaSuccess;
+}
+} // namespace o2::gpu::internal
+
+#endif

GPU/GPUTracking/Base/cuda/GPUReconstructionCUDAHelpers.inc

@@ -22,7 +22,7 @@
 #include <vector>
 #include <memory>
 #include <string>
-#include "GPUCommonChkErr.h"
+#include "GPUCommonHelpers.h"
 
 namespace o2::gpu
 {