Cleanup separate kernel file and try to create some APIs

mconcas · Sep 2, 2024 · 96c55d9 · 96c55d9
1 parent 44fe1be
commit 96c55d9
Show file tree

Hide file tree

Showing 10 changed files with 112 additions and 132 deletions.
diff --git a/Common/DCAFitter/GPU/cuda/CMakeLists.txt b/Common/DCAFitter/GPU/cuda/CMakeLists.txt
@@ -9,41 +9,43 @@
 # granted to it by virtue of its status as an Intergovernmental Organization
 # or submit itself to any jurisdiction.
 
-o2_add_library(DCAFitterNKernelsCUDA
+o2_add_library(DCAFitterCUDA
                TARGETVARNAME targetName
-               SOURCES DCAFitterNKernels.cu
-               PUBLIC_LINK_LIBRARIES O2::DCAFitter
-               PRIVATE_LINK_LIBRARIES 
-                                      O2::MathUtils
-                                      O2::ReconstructionDataFormats
-                                      O2::GPUCommon
-                                      O2::DetectorsBase
-                                      O2::GPUTrackingCUDAExternalProvider)
+               SOURCES DCAFitterN.cu
+               PUBLIC_INCLUDE_DIRECTORIES ../../include
+               PUBLIC_LINK_LIBRARIES O2::MathUtils
+                                     O2::ReconstructionDataFormats
+                                     O2::DetectorsBase
+               PRIVATE_LINK_LIBRARIES O2::GPUTrackingCUDAExternalProvider)
 set_property(TARGET ${targetName} PROPERTY CUDA_SEPARABLE_COMPILATION ON)
 
 o2_add_library(DCAFitterAPICUDA
+              # TARGETVARNAME targetName2
                SOURCES DCAFitterGPUAPI.cu
                PUBLIC_INCLUDE_DIRECTORIES ${CMAKE_SOURCE_DIR}/Common/DCAFitter/GPU
-               PUBLIC_LINK_LIBRARIES O2::GPUCommon 
+               PUBLIC_LINK_LIBRARIES O2::GPUCommon
                                      O2::ReconstructionDataFormats
-               PRIVATE_LINK_LIBRARIES O2::DCAFitterNKernelsCUDA)
+               PRIVATE_LINK_LIBRARIES O2::DCAFitterCUDA
+               # O2::GPUTrackingCUDAExternalProvider
+                )
+# set_property(TARGET ${targetName2} PROPERTY CUDA_SEPARABLE_COMPILATION ON)
 
 o2_add_test(DCAFitterNCUDA
             SOURCES test/testDCAFitterNGPU.cxx
             PUBLIC_LINK_LIBRARIES O2::ReconstructionDataFormats
                                   O2::DCAFitterAPICUDA
                                   O2::DCAFitter
                                   ROOT::Core
-                                  ROOT::Physics 
-            COMPONENT_NAME gpu               
+                                  ROOT::Physics
+            COMPONENT_NAME gpu
             LABELS vertexing
             ENVIRONMENT O2_ROOT=${CMAKE_BINARY_DIR}/stage
             VMCWORKDIR=${CMAKE_BINARY_DIR}/stage/${CMAKE_INSTALL_DATADIR})
 
 o2_add_executable(dca-fitter-on-gpu-cuda
                   SOURCES test/main_prog.cxx
                   PUBLIC_LINK_LIBRARIES O2::ReconstructionDataFormats
-                  O2::DCAFitterAPICUDA
-                  O2::DCAFitter
-                  ROOT::Core
-                  ROOT::Physics)
+                                        O2::DCAFitterAPICUDA
+                                        O2::DCAFitter
+                                        ROOT::Core
+                                        ROOT::Physics)
diff --git a/Common/DCAFitter/GPU/cuda/DCAFitterGPUAPI.cu b/Common/DCAFitter/GPU/cuda/DCAFitterGPUAPI.cu
@@ -16,7 +16,7 @@
 #endif
 
 #include "DCAFitter/DCAFitterN.h"
-#include "DCAFitterNKernels.h"
+
 #include "ReconstructionDataFormats/Track.h"
 
 #include <iostream>
@@ -36,27 +36,27 @@ inline void gpuAssert(cudaError_t code, const char* file, int line, bool abort =
   }
 }
 
-namespace o2::vertexing
+namespace o2::vertexing::gpu
 {
 /// Print fitter
-void doPrintOnGPU(o2::vertexing::DCAFitterN<2>* ft)
+void doPrintOnDevice(o2::vertexing::DCAFitterN<2>* ft)
 {
-  o2::vertexing::DCAFitterN<2>* ft_device;
+  DCAFitterN<2>* ft_device;
   gpuCheckError(cudaMalloc(reinterpret_cast<void**>(&ft_device), sizeof(o2::vertexing::DCAFitterN<2>)));
   gpuCheckError(cudaMemcpy(ft_device, ft, sizeof(o2::vertexing::DCAFitterN<2>), cudaMemcpyHostToDevice));
-  printf(" =============== GPU DCA Fitter ================\n");
-  gpu::printKernel<<<1, 1>>>(ft_device);
-  printf(" ===============================================\n");
+
+  gpu::kernel::printKernel<<<1, 1>>>(ft_device);
+
   gpuCheckError(cudaPeekAtLastError());
   gpuCheckError(cudaDeviceSynchronize());
 
   gpuCheckError(cudaFree(ft_device));
 }
 
 /// Call the process(track, ...) method
-int doProcessOnGPU(o2::vertexing::DCAFitterN<2>* ft, o2::track::TrackParCov* t1, o2::track::TrackParCov* t2)
+int doProcessOnDevice(o2::vertexing::DCAFitterN<2>* ft, o2::track::TrackParCov* t1, o2::track::TrackParCov* t2)
 {
-  o2::vertexing::DCAFitterN<2>* ft_device;
+  DCAFitterN<2>* ft_device;
   o2::track::TrackParCov* t1_device;
   o2::track::TrackParCov* t2_device;
   int result, *result_device;
@@ -70,7 +70,7 @@ int doProcessOnGPU(o2::vertexing::DCAFitterN<2>* ft, o2::track::TrackParCov* t1,
   gpuCheckError(cudaMemcpy(t1_device, t1, sizeof(o2::track::TrackParCov), cudaMemcpyHostToDevice));
   gpuCheckError(cudaMemcpy(t2_device, t2, sizeof(o2::track::TrackParCov), cudaMemcpyHostToDevice));
 
-  gpu::processKernel<<<1, 1>>>(ft_device, t1_device, t2_device, result_device);
+  gpu::kernel::processKernel<<<1, 1>>>(ft_device, t1_device, t2_device, result_device);
   gpuCheckError(cudaPeekAtLastError());
   gpuCheckError(cudaDeviceSynchronize());
 

diff --git a/Common/DCAFitter/GPU/cuda/DCAFitterGPUAPI.h b/Common/DCAFitter/GPU/cuda/DCAFitterGPUAPI.h
@@ -12,9 +12,9 @@
 #ifndef DCAFITTERN_GPU_API_H_
 #define DCAFITTERN_GPU_API_H_
 
-namespace o2::vertexing
+namespace o2::vertexing::gpu
 {
-void doPrintOnGPU(o2::vertexing::DCAFitterN<2>* ft);
-int doProcessOnGPU(o2::vertexing::DCAFitterN<2>* ft, o2::track::TrackParCov* t1, o2::track::TrackParCov* t2);
+void doPrintOnDevice(o2::vertexing::DCAFitterN<2>* ft);
+int doProcessOnDevice(o2::vertexing::DCAFitterN<2>* ft, o2::track::TrackParCov* t1, o2::track::TrackParCov* t2);
 } // namespace o2::vertexing
 #endif
diff --git a/Common/DCAFitter/GPU/cuda/DCAFitterN.cu b/Common/DCAFitter/GPU/cuda/DCAFitterN.cu
@@ -17,20 +17,38 @@
 
 #include "GPUCommonDef.h"
 #include "DCAFitter/DCAFitterN.h"
-#include "MathUtils/SMatrixGPU.h"
+// #include "MathUtils/SMatrixGPU.h"
 
-namespace o2
+#define gpuCheckError(x)                \
+  {                                     \
+    gpuAssert((x), __FILE__, __LINE__); \
+  }
+inline void gpuAssert(cudaError_t code, const char* file, int line, bool abort = true)
 {
-namespace vertexing
+  if (code != cudaSuccess) {
+    std::cout << "GPUassert: " << cudaGetErrorString(code) << " " << file << " " << line << std::endl;
+    if (abort) {
+      throw std::runtime_error("GPU assert failed.");
+    }
+  }
+}
+namespace o2::vertexing::gpu
+{
+namespace kernel
+{
+GPUg() void printKernel(o2::vertexing::DCAFitterN<2>* ft)
 {
-GPUd() void __dummy_instance__()
+  if (threadIdx.x == 0) {
+    printf(" =============== GPU DCA Fitter ================\n");
+    ft->print();
+    printf(" ===============================================\n");
+  }
+}
+
+GPUg() void processKernel(o2::vertexing::DCAFitterN<2>* ft, o2::track::TrackParCov* t1, o2::track::TrackParCov* t2, int* res)
 {
-  DCAFitter2 ft2;
-  DCAFitter3 ft3;
-  o2::track::TrackParCov tr;
-  ft2.process(tr, tr);
-  ft3.process(tr, tr, tr);
+  *res = ft->process(*t1, *t2);
 }
+} // namespace kernel
 
-} // namespace vertexing
-} // namespace o2
+} // namespace o2::vertexing::gpu
diff --git a/Common/DCAFitter/GPU/cuda/DCAFitterNKernels.cu b/Common/DCAFitter/GPU/cuda/DCAFitterNKernels.cu
diff --git a/Common/DCAFitter/GPU/cuda/DCAFitterNKernels.h b/Common/DCAFitter/GPU/cuda/DCAFitterNKernels.h
diff --git a/Common/DCAFitter/GPU/cuda/test/main_prog.cxx b/Common/DCAFitter/GPU/cuda/test/main_prog.cxx
@@ -1,3 +1,14 @@
+// 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.
+
 #include "DCAFitter/DCAFitterN.h"
 #include "CommonUtils/TreeStreamRedirector.h"
 #include <TRandom.h>
@@ -102,7 +113,7 @@ int run()
 
   double bz = 5.0;
   std::vector<int> forceQ{1, 1};
-
+  std::cout << "running... " << std::endl;
   o2::vertexing::DCAFitterN<2> ft; // 2 prong fitter
   ft.setBz(bz);
   ft.setPropagateToPCA(true);  // After finding the vertex, propagate tracks to the DCA. This is default anyway
@@ -111,14 +122,15 @@ int run()
   ft.setMaxDXYIni(4);          // do not consider V0 seeds with tracks XY-distance exceeding this. This is default anyway
   ft.setMinParamChange(1e-3);  // stop iterations if max correction is below this value. This is default anyway
   ft.setMinRelChi2Change(0.9); // stop iterations if chi2 improves by less that this factor
-
+  std::cout << "running... " << std::endl;
   auto genParent = generate(vtxGen, vctracks, bz, genPHS, k0, k0dec, forceQ);
+  std::cout << "running... " << std::endl;
   ft.setUseAbsDCA(true);
   auto res = ft.process(vctracks[0], vctracks[1]);
   ft.print();
   std::cout << "returned value: " << res << std::endl;
 
-  doPrintOnGPU(&ft);
+  o2::vertexing::gpu::doPrintOnDevice(&ft);
   return 0;
 }
 } // namespace o2::vertexing

diff --git a/Common/DCAFitter/GPU/cuda/test/testDCAFitterNGPU.cxx b/Common/DCAFitter/GPU/cuda/test/testDCAFitterNGPU.cxx
@@ -189,7 +189,7 @@ BOOST_AUTO_TEST_CASE(DCAFitterNProngs)
       auto genParent = generate(vtxGen, vctracks, bz, genPHS, k0, k0dec, forceQ);
       ft.setUseAbsDCA(true);
       swA.Start(false);
-      int ncA = doProcessOnGPU(&ft, &(vctracks[0]), &(vctracks[1]));
+      int ncA = o2::vertexing::gpu::doProcessOnDevice(&ft, &(vctracks[0]), &(vctracks[1]));
       // int ncA = ft.process(vctracks[0], vctracks[1]); // HERE WE FIT THE VERTICES
       swA.Stop();
       LOG(debug) << "fit abs.dist " << iev << " NC: " << ncA << " Chi2: " << (ncA ? ft.getChi2AtPCACandidate(0) : -1);
@@ -202,7 +202,7 @@ BOOST_AUTO_TEST_CASE(DCAFitterNProngs)
       ft.setUseAbsDCA(true);
       ft.setWeightedFinalPCA(true);
       swAW.Start(false);
-      int ncAW = doProcessOnGPU(&ft, &(vctracks[0]), &(vctracks[1]));
+      int ncAW = o2::vertexing::gpu::doProcessOnDevice(&ft, &(vctracks[0]), &(vctracks[1]));
       // int ncAW = ft.process(vctracks[0], vctracks[1]); // HERE WE FIT THE VERTICES
       swAW.Stop();
       LOG(debug) << "fit abs.dist with final weighted DCA " << iev << " NC: " << ncAW << " Chi2: " << (ncAW ? ft.getChi2AtPCACandidate(0) : -1);
@@ -215,7 +215,7 @@ BOOST_AUTO_TEST_CASE(DCAFitterNProngs)
       ft.setUseAbsDCA(false);
       ft.setWeightedFinalPCA(false);
       swW.Start(false);
-      int ncW = doProcessOnGPU(&ft, &(vctracks[0]), &(vctracks[1]));
+      int ncW = o2::vertexing::gpu::doProcessOnDevice(&ft, &(vctracks[0]), &(vctracks[1]));
       // int ncW = ft.process(vctracks[0], vctracks[1]); // HERE WE FIT THE VERTICES
       swW.Stop();
       LOG(debug) << "fit wgh.dist " << iev << " NC: " << ncW << " Chi2: " << (ncW ? ft.getChi2AtPCACandidate(0) : -1);
@@ -226,7 +226,7 @@ BOOST_AUTO_TEST_CASE(DCAFitterNProngs)
       }
     }
     ft.print();
-    doPrintOnGPU(&ft);
+    o2::vertexing::gpu::doPrintOnDevice(&ft);
     meanDA /= nfoundA ? nfoundA : 1;
     meanDAW /= nfoundA ? nfoundA : 1;
     meanDW /= nfoundW ? nfoundW : 1;

diff --git a/Common/DCAFitter/GPU/hip/CMakeLists.txt b/Common/DCAFitter/GPU/hip/CMakeLists.txt
@@ -9,29 +9,28 @@
 # granted to it by virtue of its status as an Intergovernmental Organization
 # or submit itself to any jurisdiction.
 
-# set(CMAKE_HIP_FLAGS "${CMAKE_HIP_FLAGS} -fgpu-rdc")
-# o2_add_hipified_library(DCAFitterNKernelsHIP
-#                TARGETVARNAME targetName
-#                SOURCES ../cuda/DCAFitterNKernels.cu
-#                PUBLIC_INCLUDE_DIRECTORIES ../cuda
-#                PRIVATE_LINK_LIBRARIES O2::DCAFitter
-#                                       O2::MathUtils
-#                                       O2::ReconstructionDataFormats
-#                                       O2::GPUCommon
-#                                       O2::DetectorsBase
-#                                       O2::GPUTrackingHIPExternalProvider)
-# target_compile_options(DCAFitterNKernelsHIP PRIVATE $<$<COMPILE_LANGUAGE:HIP>:-fgpu-rdc>)
-# target_link_options(DCAFitterNKernelsHIP PRIVATE $<$<COMPILE_LANGUAGE:HIP>:-fgpu-rdc>)
+set(CMAKE_HIP_FLAGS "${CMAKE_HIP_FLAGS} -fgpu-rdc")
+o2_add_hipified_library(DCAFitterHIP
+                        SOURCES ../cuda/DCAFitterN.cu
+                        PUBLIC_INCLUDE_DIRECTORIES ../../include
+                        PUBLIC_LINK_LIBRARIES O2::MathUtils
+                                              O2::ReconstructionDataFormats
+                                              O2::DetectorsBase
+                                              hip::host
+                        PRIVATE_LINK_LIBRARIES O2::GPUTrackingHIPExternalProvider
+                        TARGETVARNAME targetNAme)
+# target_compile_options(${targetName} PRIVATE $<$<COMPILE_LANGUAGE:HIP>:-fgpu-rdc>)
+# target_link_options(${targetName} PRIVATE $<$<COMPILE_LANGUAGE:HIP>:-fgpu-rdc>)
 
 # o2_add_hipified_library(DCAFitterAPIHIP
-#                SOURCES ../cuda/DCAFitterGPUAPI.cu
-#                # src/FwdDCAFitterN.cxx
-#                PUBLIC_INCLUDE_DIRECTORIES ${CMAKE_SOURCE_DIR}/Common/DCAFitter/GPU
-#                                           ../cuda
-#                PUBLIC_LINK_LIBRARIES O2::GPUCommon
-#                                      O2::ReconstructionDataFormats
-#                PRIVATE_LINK_LIBRARIES O2::DCAFitter
-#                                       O2::DCAFitterNKernelsHIP)
+#                         SOURCES ../cuda/DCAFitterGPUAPI.cu
+#                         ../cuda/DCAFitterN.cu
+#                         PUBLIC_INCLUDE_DIRECTORIES ${CMAKE_SOURCE_DIR}/Common/DCAFitter/GPU
+#                                                    ../cuda
+#                         PUBLIC_LINK_LIBRARIES O2::GPUCommon
+#                                               O2::ReconstructionDataFormats
+#                         PRIVATE_LINK_LIBRARIES O2::DCAFitterHIP O2::GPUTrackingHIPExternalProvider
+#                         TARGETVARNAME targetName)
 
 # o2_add_test(DCAFitterNHIP
 #             SOURCES ../cuda/test/testDCAFitterNGPU.cxx
@@ -41,7 +40,10 @@
 #                                   ROOT::Core
 #                                   ROOT::Physics
 #             HIPIFIED test
-#             COMPONENT_NAME gpu               
+#             COMPONENT_NAME gpu
 #             LABELS vertexing
 #             ENVIRONMENT O2_ROOT=${CMAKE_BINARY_DIR}/stage
-#             VMCWORKDIR=${CMAKE_BINARY_DIR}/stage/${CMAKE_INSTALL_DATADIR})
+#             VMCWORKDIR=${CMAKE_BINARY_DIR}/stage/${CMAKE_INSTALL_DATADIR})
+
+
+