Merge pull request #47325 from fwyzard/RecoLocalTracker_SiPixelCluste…

…rizer_GPU_DEBUG Fix `SiPixelClusterizer` alpaka code when `GPU_DEBUG` is defined
cms-sw · Feb 13, 2025 · 25fbf66 · 25fbf66
2 parents 69407d5 + 9ffd022
commit 25fbf66
Show file tree

Hide file tree

Showing 3 changed files with 41 additions and 13 deletions.
diff --git a/FWCore/Utilities/interface/DeviceGlobal.h b/FWCore/Utilities/interface/DeviceGlobal.h
@@ -0,0 +1,28 @@
+#ifndef FWCore_Utilities_DeviceGlobal_h
+#define FWCore_Utilities_DeviceGlobal_h
+
+// FIXME alpaka provides ALPAKA_STATIC_ACC_MEM_GLOBAL to declare device global
+// variables, but it is currently not working as expected. Improve its behaviour
+// and syntax and migrate to that.
+
+#if defined(__SYCL_DEVICE_ONLY__)
+
+// The SYCL standard does not support device global variables.
+// oneAPI defines the sycl_ext_oneapi_device_global extension, but with an awkward syntax
+// that is not easily compatible with CUDA, HIP and regular C++ global variables.
+#error "The SYCL backend does not support device global variables"
+#define DEVICE_GLOBAL
+
+#elif defined(__CUDA_ARCH__) or defined(__HIP_DEVICE_COMPILE__)
+
+// CUDA and HIP/ROCm device compilers use the __device__ attribute.
+#define DEVICE_GLOBAL __device__
+
+#else
+
+// host compilers do not need any special attributes.
+#define DEVICE_GLOBAL
+
+#endif
+
+#endif  // FWCore_Utilities_DeviceGlobal_h
diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/PixelClustering.h b/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/PixelClustering.h
@@ -9,6 +9,7 @@
 #include <alpaka/alpaka.hpp>
 
 #include "DataFormats/SiPixelClusterSoA/interface/ClusteringConstants.h"
+#include "FWCore/Utilities/interface/DeviceGlobal.h"
 #include "FWCore/Utilities/interface/HostDeviceConstant.h"
 #include "Geometry/CommonTopologies/interface/SimplePixelTopology.h"
 #include "HeterogeneousCore/AlpakaInterface/interface/HistoContainer.h"
@@ -20,7 +21,7 @@
 namespace ALPAKA_ACCELERATOR_NAMESPACE::pixelClustering {
 
 #ifdef GPU_DEBUG
-  HOST_DEVICE_CONSTANT uint32_t gMaxHit = 0;
+  DEVICE_GLOBAL uint32_t gMaxHit = 0;
 #endif
 
   namespace pixelStatus {

diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/SiPixelRawToClusterKernel.dev.cc b/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/SiPixelRawToClusterKernel.dev.cc
@@ -35,7 +35,7 @@
 #include "PixelClustering.h"
 #include "SiPixelRawToClusterKernel.h"
 
-// #define GPU_DEBUG
+//#define GPU_DEBUG
 
 namespace ALPAKA_ACCELERATOR_NAMESPACE {
   namespace pixelDetails {
@@ -493,21 +493,20 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE {
           alpaka::syncBlockThreads(acc);
         }
 #ifdef GPU_DEBUG
-        ALPAKA_ASSERT_ACC(0 == clus_view[1].moduleStart());
-        auto c0 = std::min(maxHitsInModule, clus_view[2].clusModuleStart());
-        ALPAKA_ASSERT_ACC(c0 == clus_view[2].moduleStart());
-        ALPAKA_ASSERT_ACC(clus_view[1024].moduleStart() >= clus_view[1023].moduleStart());
-        ALPAKA_ASSERT_ACC(clus_view[1025].moduleStart() >= clus_view[1024].moduleStart());
-        ALPAKA_ASSERT_ACC(clus_view[numberOfModules].moduleStart() >= clus_view[1025].moduleStart());
-
-        for (uint32_t i : cms::alpakatools::independent_group_elements(acc, numberOfModules + 1)) {
-          if (0 != i)
-            ALPAKA_ASSERT_ACC(clus_view[i].moduleStart() >= clus_view[i - 1].moduleStart());
+        ALPAKA_ASSERT_ACC(0 == clus_view[0].clusModuleStart());
+        auto c0 = std::min(maxHitsInModule, clus_view[1].clusModuleStart());
+        ALPAKA_ASSERT_ACC(c0 == clus_view[1].clusModuleStart());
+        ALPAKA_ASSERT_ACC(clus_view[1024].clusModuleStart() >= clus_view[1023].clusModuleStart());
+        ALPAKA_ASSERT_ACC(clus_view[1025].clusModuleStart() >= clus_view[1024].clusModuleStart());
+        ALPAKA_ASSERT_ACC(clus_view[numberOfModules].clusModuleStart() >= clus_view[1025].clusModuleStart());
+
+        for (uint32_t i : cms::alpakatools::independent_group_elements(acc, numberOfModules)) {
+          ALPAKA_ASSERT_ACC(clus_view[i + 1].clusModuleStart() >= clus_view[i].clusModuleStart());
           // Check BPX2 (1), FP1 (4)
           constexpr auto bpix2 = TrackerTraits::layerStart[1];
           constexpr auto fpix1 = TrackerTraits::layerStart[4];
           if (i == bpix2 || i == fpix1)
-            printf("moduleStart %d %d\n", i, clus_view[i].moduleStart());
+            printf("moduleStart %d %d\n", i, clus_view[i].clusModuleStart());
         }
 
 #endif