GPU OpenCL: Reorganize OpenCL code and remove obsolete internals header

Author: davidrohr

GPU/GPUTracking/Base/opencl/CMakeLists.txt

@@ -37,7 +37,7 @@ set(OCL_DEFINECL "-D$<JOIN:$<TARGET_PROPERTY:O2::GPUTracking,COMPILE_DEFINITIONS
 )
 
 set(SRCS GPUReconstructionOCL.cxx)
-set(HDRS GPUReconstructionOCL.h GPUReconstructionOCLInternals.h)
+set(HDRS GPUReconstructionOCL.h)
 
 if (NOT DEFINED GPUCA_OCL_SPIRV_VERSION)
   set(GPUCA_OCL_SPIRV_VERSION 1.2)

GPU/GPUTracking/Base/opencl/GPUReconstructionOCL.cxx

@@ -15,8 +15,15 @@
 #define GPUCA_GPUTYPE_OPENCL
 #define __OPENCL_HOST__
 
+#define CL_TARGET_OPENCL_VERSION 220
+#include <CL/opencl.h>
+#include <CL/cl_ext.h>
+#include <vector>
+#include <string>
+#include <memory>
+#include "GPULogging.h"
+
 #include "GPUReconstructionOCL.h"
-#include "GPUReconstructionOCLInternals.h"
 #include "GPUReconstructionIncludes.h"
 
 using namespace o2::gpu;
@@ -26,6 +33,154 @@ using namespace o2::gpu;
 #include <typeinfo>
 #include <cstdlib>
 
+namespace o2::gpu
+{
+
+static const char* opencl_error_string(int32_t errorcode)
+{
+  switch (errorcode) {
+    case CL_SUCCESS:
+      return "Success!";
+    case CL_DEVICE_NOT_FOUND:
+      return "Device not found.";
+    case CL_DEVICE_NOT_AVAILABLE:
+      return "Device not available";
+    case CL_COMPILER_NOT_AVAILABLE:
+      return "Compiler not available";
+    case CL_MEM_OBJECT_ALLOCATION_FAILURE:
+      return "Memory object allocation failure";
+    case CL_OUT_OF_RESOURCES:
+      return "Out of resources";
+    case CL_OUT_OF_HOST_MEMORY:
+      return "Out of host memory";
+    case CL_PROFILING_INFO_NOT_AVAILABLE:
+      return "Profiling information not available";
+    case CL_MEM_COPY_OVERLAP:
+      return "Memory copy overlap";
+    case CL_IMAGE_FORMAT_MISMATCH:
+      return "Image format mismatch";
+    case CL_IMAGE_FORMAT_NOT_SUPPORTED:
+      return "Image format not supported";
+    case CL_BUILD_PROGRAM_FAILURE:
+      return "Program build failure";
+    case CL_MAP_FAILURE:
+      return "Map failure";
+    case CL_INVALID_VALUE:
+      return "Invalid value";
+    case CL_INVALID_DEVICE_TYPE:
+      return "Invalid device type";
+    case CL_INVALID_PLATFORM:
+      return "Invalid platform";
+    case CL_INVALID_DEVICE:
+      return "Invalid device";
+    case CL_INVALID_CONTEXT:
+      return "Invalid context";
+    case CL_INVALID_QUEUE_PROPERTIES:
+      return "Invalid queue properties";
+    case CL_INVALID_COMMAND_QUEUE:
+      return "Invalid command queue";
+    case CL_INVALID_HOST_PTR:
+      return "Invalid host pointer";
+    case CL_INVALID_MEM_OBJECT:
+      return "Invalid memory object";
+    case CL_INVALID_IMAGE_FORMAT_DESCRIPTOR:
+      return "Invalid image format descriptor";
+    case CL_INVALID_IMAGE_SIZE:
+      return "Invalid image size";
+    case CL_INVALID_SAMPLER:
+      return "Invalid sampler";
+    case CL_INVALID_BINARY:
+      return "Invalid binary";
+    case CL_INVALID_BUILD_OPTIONS:
+      return "Invalid build options";
+    case CL_INVALID_PROGRAM:
+      return "Invalid program";
+    case CL_INVALID_PROGRAM_EXECUTABLE:
+      return "Invalid program executable";
+    case CL_INVALID_KERNEL_NAME:
+      return "Invalid kernel name";
+    case CL_INVALID_KERNEL_DEFINITION:
+      return "Invalid kernel definition";
+    case CL_INVALID_KERNEL:
+      return "Invalid kernel";
+    case CL_INVALID_ARG_INDEX:
+      return "Invalid argument index";
+    case CL_INVALID_ARG_VALUE:
+      return "Invalid argument value";
+    case CL_INVALID_ARG_SIZE:
+      return "Invalid argument size";
+    case CL_INVALID_KERNEL_ARGS:
+      return "Invalid kernel arguments";
+    case CL_INVALID_WORK_DIMENSION:
+      return "Invalid work dimension";
+    case CL_INVALID_WORK_GROUP_SIZE:
+      return "Invalid work group size";
+    case CL_INVALID_WORK_ITEM_SIZE:
+      return "Invalid work item size";
+    case CL_INVALID_GLOBAL_OFFSET:
+      return "Invalid global offset";
+    case CL_INVALID_EVENT_WAIT_LIST:
+      return "Invalid event wait list";
+    case CL_INVALID_EVENT:
+      return "Invalid event";
+    case CL_INVALID_OPERATION:
+      return "Invalid operation";
+    case CL_INVALID_GL_OBJECT:
+      return "Invalid OpenGL object";
+    case CL_INVALID_BUFFER_SIZE:
+      return "Invalid buffer size";
+    case CL_INVALID_MIP_LEVEL:
+      return "Invalid mip-map level";
+    default:
+      return "Unknown Errorcode";
+  }
+}
+
+#define GPUFailedMsg(x) GPUFailedMsgA(x, __FILE__, __LINE__)
+#define GPUFailedMsgI(x) GPUFailedMsgAI(x, __FILE__, __LINE__)
+
+static inline int64_t OCLsetKernelParameters_helper(cl_kernel& k, int32_t i)
+{
+  return 0;
+}
+
+template <typename T, typename... Args>
+static inline int64_t OCLsetKernelParameters_helper(cl_kernel& kernel, int32_t i, const T& firstParameter, const Args&... restOfParameters)
+{
+  int64_t retVal = clSetKernelArg(kernel, i, sizeof(T), &firstParameter);
+  if (retVal) {
+    return retVal;
+  }
+  return OCLsetKernelParameters_helper(kernel, i + 1, restOfParameters...);
+}
+
+template <typename... Args>
+static inline int64_t OCLsetKernelParameters(cl_kernel& kernel, const Args&... args)
+{
+  return OCLsetKernelParameters_helper(kernel, 0, args...);
+}
+
+static inline int64_t clExecuteKernelA(cl_command_queue queue, cl_kernel krnl, size_t local_size, size_t global_size, cl_event* pEvent, cl_event* wait = nullptr, cl_int nWaitEvents = 1)
+{
+  return clEnqueueNDRangeKernel(queue, krnl, 1, nullptr, &global_size, &local_size, wait == nullptr ? 0 : nWaitEvents, wait, pEvent);
+}
+
+struct GPUReconstructionOCLInternals {
+  cl_platform_id platform;
+  cl_device_id device;
+  cl_context context;
+  cl_command_queue command_queue[GPUCA_MAX_STREAMS];
+  cl_mem mem_gpu;
+  cl_mem mem_constant;
+  cl_mem mem_host;
+  cl_program program;
+
+  std::vector<std::pair<cl_kernel, std::string>> kernels;
+};
+
+static_assert(std::is_convertible<cl_event, void*>::value, "OpenCL event type incompatible to deviceEvent");
+} // namespace o2::gpu
+
 #define GPUErrorReturn(...) \
   {                         \
     GPUError(__VA_ARGS__);  \
@@ -43,6 +198,77 @@ QGET_LD_BINARY_SYMBOLS(GPUReconstructionOCLCode_src);
 QGET_LD_BINARY_SYMBOLS(GPUReconstructionOCLCode_spirv);
 #endif
 
+template <class T, int32_t I, typename... Args>
+inline void GPUReconstructionOCLBackend::runKernelBackendInternal(const krnlSetupTime& _xyz, const Args&... args)
+{
+  cl_kernel k = _xyz.y.num > 1 ? getKernelObject<cl_kernel, T, I, true>() : getKernelObject<cl_kernel, T, I, false>();
+  auto& x = _xyz.x;
+  auto& y = _xyz.y;
+  auto& z = _xyz.z;
+  if (y.num <= 1) {
+    GPUFailedMsg(OCLsetKernelParameters(k, mInternals->mem_gpu, mInternals->mem_constant, y.start, args...));
+  } else {
+    GPUFailedMsg(OCLsetKernelParameters(k, mInternals->mem_gpu, mInternals->mem_constant, y.start, y.num, args...));
+  }
+
+  cl_event ev;
+  cl_event* evr;
+  bool tmpEvent = false;
+  if (z.ev == nullptr && mProcessingSettings.deviceTimers && mProcessingSettings.debugLevel > 0) {
+    evr = &ev;
+    tmpEvent = true;
+  } else {
+    evr = (cl_event*)z.ev;
+  }
+  GPUFailedMsg(clExecuteKernelA(mInternals->command_queue[x.stream], k, x.nThreads, x.nThreads * x.nBlocks, evr, (cl_event*)z.evList, z.nEvents));
+  if (mProcessingSettings.deviceTimers && mProcessingSettings.debugLevel > 0) {
+    cl_ulong time_start, time_end;
+    GPUFailedMsg(clWaitForEvents(1, evr));
+    GPUFailedMsg(clGetEventProfilingInfo(*evr, CL_PROFILING_COMMAND_START, sizeof(time_start), &time_start, nullptr));
+    GPUFailedMsg(clGetEventProfilingInfo(*evr, CL_PROFILING_COMMAND_END, sizeof(time_end), &time_end, nullptr));
+    _xyz.t = (time_end - time_start) * 1.e-9f;
+    if (tmpEvent) {
+      GPUFailedMsg(clReleaseEvent(ev));
+    }
+  }
+}
+
+template <class T, int32_t I>
+int32_t GPUReconstructionOCLBackend::AddKernel(bool multi)
+{
+  std::string name(GetKernelName<T, I>());
+  if (multi) {
+    name += "_multi";
+  }
+  std::string kname("krnl_" + name);
+
+  cl_int ocl_error;
+  cl_kernel krnl = clCreateKernel(mInternals->program, kname.c_str(), &ocl_error);
+  if (GPUFailedMsgI(ocl_error)) {
+    GPUError("Error creating OPENCL Kernel: %s", name.c_str());
+    return 1;
+  }
+  mInternals->kernels.emplace_back(krnl, name);
+  return 0;
+}
+
+template <class T, int32_t I>
+inline uint32_t GPUReconstructionOCLBackend::FindKernel(int32_t num)
+{
+  std::string name(GetKernelName<T, I>());
+  if (num > 1) {
+    name += "_multi";
+  }
+
+  for (uint32_t k = 0; k < mInternals->kernels.size(); k++) {
+    if (mInternals->kernels[k].second == name) {
+      return (k);
+    }
+  }
+  GPUError("Could not find OpenCL kernel %s", name.c_str());
+  throw ::std::runtime_error("Requested unsupported OpenCL kernel");
+}
+
 GPUReconstruction* GPUReconstruction_Create_OCL(const GPUSettingsDeviceBackend& cfg) { return new GPUReconstructionOCL(cfg); }
 
 GPUReconstructionOCLBackend::GPUReconstructionOCLBackend(const GPUSettingsDeviceBackend& cfg) : GPUReconstructionDeviceBase(cfg, sizeof(GPUReconstructionDeviceBase))
@@ -561,8 +787,7 @@ int32_t GPUReconstructionOCLBackend::GPUDebug(const char* state, int32_t stream,
 template <class T, int32_t I, typename... Args>
 void GPUReconstructionOCLBackend::runKernelBackend(const krnlSetupArgs<T, I, Args...>& args)
 {
-  cl_kernel k = args.s.y.num > 1 ? getKernelObject<cl_kernel, T, I, true>() : getKernelObject<cl_kernel, T, I, false>();
-  std::apply([this, &args, &k](auto&... vals) { runKernelBackendInternal(args.s, k, vals...); }, args.v);
+  std::apply([this, &args](auto&... vals) { runKernelBackendInternal<T, I, Args...>(args.s, vals...); }, args.v);
 }
 
 template <class S, class T, int32_t I, bool MULTI>

GPU/GPUTracking/Base/opencl/GPUReconstructionOCL.h

@@ -59,8 +59,8 @@ class GPUReconstructionOCLBackend : public GPUReconstructionDeviceBase
   int32_t AddKernel(bool multi = false);
   template <class T, int32_t I = 0>
   uint32_t FindKernel(int32_t num);
-  template <typename K, typename... Args>
-  void runKernelBackendInternal(const krnlSetupTime& _xyz, K& k, const Args&... args);
+  template <class T, int32_t I = 0, typename... Args>
+  void runKernelBackendInternal(const krnlSetupTime& _xyz, const Args&... args);
   template <class T, int32_t I = 0>
   gpu_reconstruction_kernels::krnlProperties getKernelPropertiesBackend();