Merge pull request cms-sw#29 from fwyzard/fix_setHitsLayerStart

Fix the thread loop in `setHitsLayerStart`
missirol · Jun 30, 2023 · 0b7c738 · 0b7c738
2 parents 8252595 + 687ad23
commit 0b7c738
Show file tree

Hide file tree

Showing 3 changed files with 9 additions and 12 deletions.
diff --git a/HLTrigger/Configuration/test/testHLTWithAlpakaPixelReco.sh b/HLTrigger/Configuration/test/testHLTWithAlpakaPixelReco.sh
@@ -8,7 +8,7 @@ hltGetConfiguration /frozen/2023/2e34/v1.2/HLT \
   --output all \
   --max-events 100 \
   --paths DQM_PixelReco*,*DQMGPUvsCPU* \
-  --input file:/cmsnfsgpu_data/gpu_data/store/data/Run2023C/EphemeralHLTPhysics0/RAW/v1/000/368/822/00000/6e1268da-f96a-49f6-a5f0-89933142dd89.root \
+  --input /store/data/Run2023C/EphemeralHLTPhysics0/RAW/v1/000/368/822/00000/6e1268da-f96a-49f6-a5f0-89933142dd89.root \
   --customise \
 HLTrigger/Configuration/customizeHLTforPatatrack.customiseHLTforAlpakaPixelReco,\
 HLTrigger/Configuration/customizeHLTforPatatrack.customiseHLTforTestingDQMGPUvsCPUPixelOnlyUpToLocal \

diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/alpaka/PixelRecHitGPUKernel.dev.cc b/RecoLocalTracker/SiPixelRecHits/plugins/alpaka/PixelRecHitGPUKernel.dev.cc
@@ -22,20 +22,17 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE {
     template <typename TAcc, typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
     ALPAKA_FN_ACC void operator()(TAcc const& acc,
                                   uint32_t const* __restrict__ hitsModuleStart,
-                                  pixelCPEforDevice::ParamsOnDeviceT<TrackerTraits> const* cpeParams,
-                                  uint32_t* hitsLayerStart) const {
-      const int32_t i = alpaka::getIdx<alpaka::Grid, alpaka::Threads>(acc)[0u];
-      constexpr auto m = TrackerTraits::numberOfLayers;
-
+                                  pixelCPEforDevice::ParamsOnDeviceT<TrackerTraits> const* __restrict__ cpeParams,
+                                  uint32_t* __restrict__ hitsLayerStart) const {
       assert(0 == hitsModuleStart[0]);
 
-      if (i <= (int32_t)m) {
+      for (int32_t i : cms::alpakatools::elements_with_stride(acc, TrackerTraits::numberOfLayers + 1)) {
         hitsLayerStart[i] = hitsModuleStart[cpeParams->layerGeometry().layerStart[i]];
 #ifdef GPU_DEBUG
         int old = i == 0 ? 0 : hitsModuleStart[cpeParams->layerGeometry().layerStart[i - 1]];
         printf("LayerStart %d/%d at module %d: %d - %d\n",
                i,
-               m,
+               TrackerTraits::numberOfLayers,
                cpeParams->layerGeometry().layerStart[i],
                hitsLayerStart[i],
                hitsLayerStart[i] - old);
@@ -67,7 +64,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE {
         const auto workDiv1D = cms::alpakatools::make_workdiv<Acc1D>(blocks, threadsPerBlock);
 
 #ifdef GPU_DEBUG
-        std::cout << "launching getHits kernel for " << blocks << " blocks" << std::endl;
+        std::cout << "launching getHits kernel on " << alpaka::core::demangled<Acc1D> << " with " << blocks << " blocks" << std::endl;
 #endif
         alpaka::exec<Acc1D>(queue,
                             workDiv1D,

diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/alpaka/pixelRecHits.h b/RecoLocalTracker/SiPixelRecHits/plugins/alpaka/pixelRecHits.h
@@ -80,15 +80,15 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE {
         if (0 == nclus)
           return;
 #ifdef GPU_DEBUG
-        if (threadIdx == 0) {
-          auto k = clusters[1 + blockIdx.x].moduleStart();
+        if (threadIdxLocal == 0) {
+          auto k = clusters[1 + blockIdx].moduleStart();
           while (digis[k].moduleId() == invalidModuleId)
             ++k;
           ALPAKA_ASSERT_OFFLOAD(digis[k].moduleId() == me);
         }
 
         if (me % 100 == 1)
-          if (threadIdx == 0)
+          if (threadIdxLocal == 0)
             printf(
                 "hitbuilder: %d clusters in module %d. will write at %d\n", nclus, me, clusters[me].clusModuleStart());
 #endif