From 44827dfedde8de7ad7c5d6ac8d14d8dedb72924e Mon Sep 17 00:00:00 2001 From: Andrea Bocci Date: Thu, 29 Jun 2023 22:44:33 +0200 Subject: [PATCH 1/3] Use the logical file name for the input sample --- HLTrigger/Configuration/test/testHLTWithAlpakaPixelReco.sh | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/HLTrigger/Configuration/test/testHLTWithAlpakaPixelReco.sh b/HLTrigger/Configuration/test/testHLTWithAlpakaPixelReco.sh index 78ac2537a036f..208566e79d476 100755 --- 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 \ From 0dd1511c234bfead6294e841f58d1a2706348ff3 Mon Sep 17 00:00:00 2001 From: Andrea Bocci Date: Thu, 29 Jun 2023 23:20:42 +0200 Subject: [PATCH 2/3] Fix debug statements --- .../plugins/alpaka/PixelRecHitGPUKernel.dev.cc | 2 +- .../SiPixelRecHits/plugins/alpaka/pixelRecHits.h | 6 +++--- 2 files changed, 4 insertions(+), 4 deletions(-) diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/alpaka/PixelRecHitGPUKernel.dev.cc b/RecoLocalTracker/SiPixelRecHits/plugins/alpaka/PixelRecHitGPUKernel.dev.cc index 58ade6e21a21d..56dd62df21ce5 100644 --- a/RecoLocalTracker/SiPixelRecHits/plugins/alpaka/PixelRecHitGPUKernel.dev.cc +++ b/RecoLocalTracker/SiPixelRecHits/plugins/alpaka/PixelRecHitGPUKernel.dev.cc @@ -67,7 +67,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE { const auto workDiv1D = cms::alpakatools::make_workdiv(blocks, threadsPerBlock); #ifdef GPU_DEBUG - std::cout << "launching getHits kernel for " << blocks << " blocks" << std::endl; + std::cout << "launching getHits kernel on " << alpaka::core::demangled << " with " << blocks << " blocks" << std::endl; #endif alpaka::exec(queue, workDiv1D, diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/alpaka/pixelRecHits.h b/RecoLocalTracker/SiPixelRecHits/plugins/alpaka/pixelRecHits.h index e676447934b2e..04a3893a4699d 100644 --- 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 From 687ad235944d17a2be15ebc1194339e108216b8b Mon Sep 17 00:00:00 2001 From: Andrea Bocci Date: Thu, 29 Jun 2023 23:23:32 +0200 Subject: [PATCH 3/3] Fix the thread loop in setHitsLayerStart --- .../plugins/alpaka/PixelRecHitGPUKernel.dev.cc | 11 ++++------- 1 file changed, 4 insertions(+), 7 deletions(-) diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/alpaka/PixelRecHitGPUKernel.dev.cc b/RecoLocalTracker/SiPixelRecHits/plugins/alpaka/PixelRecHitGPUKernel.dev.cc index 56dd62df21ce5..4b6f78a634c12 100644 --- 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 >> ALPAKA_FN_ACC void operator()(TAcc const& acc, uint32_t const* __restrict__ hitsModuleStart, - pixelCPEforDevice::ParamsOnDeviceT const* cpeParams, - uint32_t* hitsLayerStart) const { - const int32_t i = alpaka::getIdx(acc)[0u]; - constexpr auto m = TrackerTraits::numberOfLayers; - + pixelCPEforDevice::ParamsOnDeviceT 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);