From 9d5a7317dd960b72d0d8df3d6e1e37c287413d88 Mon Sep 17 00:00:00 2001 From: David Rohr Date: Tue, 2 Mar 2021 09:39:24 +0100 Subject: [PATCH 01/20] GPU: Improve debug output during event dump loading --- GPU/GPUTracking/Benchmark/standalone.cxx | 8 ++++++-- 1 file changed, 6 insertions(+), 2 deletions(-) diff --git a/GPU/GPUTracking/Benchmark/standalone.cxx b/GPU/GPUTracking/Benchmark/standalone.cxx index 3e1dc69bb91ed..0fe173e22c5a9 100644 --- a/GPU/GPUTracking/Benchmark/standalone.cxx +++ b/GPU/GPUTracking/Benchmark/standalone.cxx @@ -795,11 +795,15 @@ int main(int argc, char** argv) ioPtrEvents.resize(configStandalone.preloadEvents ? (nEvents - configStandalone.StartEvent) : 1); ioMemEvents.resize(configStandalone.preloadEvents ? (nEvents - configStandalone.StartEvent) : 1); if (configStandalone.preloadEvents) { - printf("Preloading events"); + printf("Preloading events%s", configStandalone.proc.debugLevel >= 2 ? "\n" : ""); fflush(stdout); for (int i = 0; i < nEvents - configStandalone.StartEvent; i++) { LoadEvent(configStandalone.StartEvent + i, i); - printf(" %d", i); + if (configStandalone.proc.debugLevel >= 2) { + printf("Loading event %d\n", i); + } else { + printf(" %d", i); + } fflush(stdout); } printf("\n"); From a35e2f4fad0316ef0e849a288fddd07f43dac2f2 Mon Sep 17 00:00:00 2001 From: David Rohr Date: Tue, 2 Mar 2021 09:41:22 +0100 Subject: [PATCH 02/20] GPU: Move configuration file to a more appropriate place --- GPU/GPUTracking/CMakeLists.txt | 4 ++-- GPU/GPUTracking/{DataTypes => Definitions}/GPUSettingsList.h | 0 2 files changed, 2 insertions(+), 2 deletions(-) rename GPU/GPUTracking/{DataTypes => Definitions}/GPUSettingsList.h (100%) diff --git a/GPU/GPUTracking/CMakeLists.txt b/GPU/GPUTracking/CMakeLists.txt index 35cb159d005a0..c70bb0de3ec46 100644 --- a/GPU/GPUTracking/CMakeLists.txt +++ b/GPU/GPUTracking/CMakeLists.txt @@ -124,7 +124,7 @@ set(HDRS_INSTALL TPCConvert/GPUTPCConvertImpl.h Base/GPUReconstructionKernelMacros.h DataTypes/GPUO2FakeClasses.h - DataTypes/GPUSettingsList.h + Definitions/GPUSettingsList.h DataTypes/GPUOutputControl.h DataTypes/GPUO2DataTypes.h DataTypes/GPUHostDataTypes.h @@ -150,7 +150,7 @@ set(HDRS_INSTALL if(ALIGPU_BUILD_TYPE STREQUAL "O2") set(SRCS ${SRCS} Interface/GPUO2InterfaceConfigurableParam.cxx) set(HDRS_CINT_O2 ${HDRS_CINT_O2} DataTypes/TPCdEdxCalibrationSplines.h Interface/GPUO2InterfaceConfigurableParam.h) - set(HDRS_CINT_O2_ADDITIONAL DataTypes/GPUSettings.h DataTypes/GPUSettingsList.h DataTypes/GPUDataTypes.h TRDTracking/GPUTRDTrack.h) # Manual dependencies for ROOT dictionary generation + set(HDRS_CINT_O2_ADDITIONAL DataTypes/GPUSettings.h Definitions/GPUSettingsList.h DataTypes/GPUDataTypes.h TRDTracking/GPUTRDTrack.h) # Manual dependencies for ROOT dictionary generation endif() # Sources for O2 and for Standalone if requested in config file diff --git a/GPU/GPUTracking/DataTypes/GPUSettingsList.h b/GPU/GPUTracking/Definitions/GPUSettingsList.h similarity index 100% rename from GPU/GPUTracking/DataTypes/GPUSettingsList.h rename to GPU/GPUTracking/Definitions/GPUSettingsList.h From e5ae644c3a8e88bfe1c2f8379d0141f274293224 Mon Sep 17 00:00:00 2001 From: David Rohr Date: Tue, 2 Mar 2021 10:37:53 +0100 Subject: [PATCH 03/20] GPU: Add some more functionality to memory size scalig class --- GPU/GPUTracking/Base/GPUReconstruction.cxx | 7 ++++ GPU/GPUTracking/CMakeLists.txt | 2 +- .../DataTypes/GPUMemorySizeScalers.cxx | 32 +++++++++++++++++++ .../DataTypes/GPUMemorySizeScalers.h | 28 ++++++++++------ GPU/GPUTracking/Definitions/GPUSettingsList.h | 1 + 5 files changed, 59 insertions(+), 11 deletions(-) create mode 100644 GPU/GPUTracking/DataTypes/GPUMemorySizeScalers.cxx diff --git a/GPU/GPUTracking/Base/GPUReconstruction.cxx b/GPU/GPUTracking/Base/GPUReconstruction.cxx index 34540d170b3bc..7bfb66b55f419 100644 --- a/GPU/GPUTracking/Base/GPUReconstruction.cxx +++ b/GPU/GPUTracking/Base/GPUReconstruction.cxx @@ -278,6 +278,10 @@ int GPUReconstruction::InitPhaseBeforeDevice() } mMemoryScalers->factor = mProcessingSettings.memoryScalingFactor; + mMemoryScalers->returnMaxVal = mProcessingSettings.forceMaxMemScalers; + if (mProcessingSettings.forceMaxMemScalers > 1) { + mMemoryScalers->rescaleMaxMem(mProcessingSettings.forceMaxMemScalers); + } #ifdef WITH_OPENMP if (mProcessingSettings.ompThreads <= 0) { @@ -347,6 +351,9 @@ int GPUReconstruction::InitPhasePermanentMemory() int GPUReconstruction::InitPhaseAfterDevice() { + if (mProcessingSettings.forceMaxMemScalers <= 1 && mProcessingSettings.memoryAllocationStrategy == GPUMemoryResource::ALLOCATION_GLOBAL) { + mMemoryScalers->rescaleMaxMem(IsGPU() ? mDeviceMemorySize : mHostMemorySize); + } for (unsigned int i = 0; i < mChains.size(); i++) { if (mChains[i]->Init()) { return 1; diff --git a/GPU/GPUTracking/CMakeLists.txt b/GPU/GPUTracking/CMakeLists.txt index c70bb0de3ec46..c3f8c5a7c2f5b 100644 --- a/GPU/GPUTracking/CMakeLists.txt +++ b/GPU/GPUTracking/CMakeLists.txt @@ -63,6 +63,7 @@ set(SRCS set(SRCS_NO_CINT DataTypes/GPUDataTypes.cxx + DataTypes/GPUMemorySizeScalers.cxx Base/GPUReconstruction.cxx Base/GPUReconstructionCPU.cxx Base/GPUProcessor.cxx @@ -128,7 +129,6 @@ set(HDRS_INSTALL DataTypes/GPUOutputControl.h DataTypes/GPUO2DataTypes.h DataTypes/GPUHostDataTypes.h - DataTypes/GPUMemorySizeScalers.h DataTypes/GPURawData.h DataTypes/GPUdEdxInfo.h Base/GPUParam.inc diff --git a/GPU/GPUTracking/DataTypes/GPUMemorySizeScalers.cxx b/GPU/GPUTracking/DataTypes/GPUMemorySizeScalers.cxx new file mode 100644 index 0000000000000..95edfdba17841 --- /dev/null +++ b/GPU/GPUTracking/DataTypes/GPUMemorySizeScalers.cxx @@ -0,0 +1,32 @@ +// Copyright CERN and copyright holders of ALICE O2. This software is +// distributed under the terms of the GNU General Public License v3 (GPL +// Version 3), copied verbatim in the file "COPYING". +// +// See http://alice-o2.web.cern.ch/license for full licensing information. +// +// 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. + +/// \file GPUMemorySizeScalers.cxx +/// \author David Rohr + +#include "GPUMemorySizeScalers.h" + +using namespace GPUCA_NAMESPACE::gpu; + +void GPUMemorySizeScalers::rescaleMaxMem(size_t newAvailableMemory) +{ + GPUMemorySizeScalers tmp; + tpcMaxPeaks = (double)tmp.tpcMaxPeaks * newAvailableMemory / tmp.availableMemory; + tpcMaxClusters = (double)tmp.tpcMaxClusters * newAvailableMemory / tmp.availableMemory; + tpcMaxStartHits = (double)tmp.tpcMaxStartHits * newAvailableMemory / tmp.availableMemory; + tpcMaxRowStartHits = (double)tmp.tpcMaxRowStartHits * newAvailableMemory / tmp.availableMemory; + tpcMaxTracklets = (double)tmp.tpcMaxTracklets * newAvailableMemory / tmp.availableMemory; + tpcMaxTrackletHits = (double)tmp.tpcMaxTrackletHits * newAvailableMemory / tmp.availableMemory; + tpcMaxSectorTracks = (double)tmp.tpcMaxSectorTracks * newAvailableMemory / tmp.availableMemory; + tpcMaxSectorTrackHits = (double)tmp.tpcMaxSectorTrackHits * newAvailableMemory / tmp.availableMemory; + tpcMaxMergedTracks = (double)tmp.tpcMaxMergedTracks * newAvailableMemory / tmp.availableMemory; + tpcMaxMergedTrackHits = (double)tmp.tpcMaxMergedTrackHits * newAvailableMemory / tmp.availableMemory; + availableMemory = newAvailableMemory; +} diff --git a/GPU/GPUTracking/DataTypes/GPUMemorySizeScalers.h b/GPU/GPUTracking/DataTypes/GPUMemorySizeScalers.h index 24f7d6b97de6b..7e1d3e786ceae 100644 --- a/GPU/GPUTracking/DataTypes/GPUMemorySizeScalers.h +++ b/GPU/GPUTracking/DataTypes/GPUMemorySizeScalers.h @@ -55,17 +55,25 @@ struct GPUMemorySizeScalers { size_t tpcMaxSectorTrackHits = 11500000; size_t tpcMaxMergedTracks = 5800000; size_t tpcMaxMergedTrackHits = 380000000; + size_t availableMemory = 24000000000; + bool returnMaxVal = false; - size_t NTPCPeaks(size_t tpcDigits) { return std::min(tpcMaxPeaks, hitOffset + tpcDigits * tpcPeaksPerDigit) * factor; } - size_t NTPCClusters(size_t tpcDigits) { return std::min(tpcMaxClusters, tpcClustersPerPeak * NTPCPeaks(tpcDigits)) * factor; } - size_t NTPCStartHits(size_t tpcHits) { return std::min(tpcMaxStartHits, offset + tpcHits * tpcStartHitsPerHit) * factor; } - size_t NTPCRowStartHits(size_t tpcHits) { return std::min(tpcMaxRowStartHits, offset + NTPCStartHits(tpcHits) / GPUCA_ROW_COUNT * 4.) * factor; } - size_t NTPCTracklets(size_t tpcHits) { return std::min(tpcMaxTracklets, NTPCStartHits(tpcHits) * tpcTrackletsPerStartHit) * factor; } - size_t NTPCTrackletHits(size_t tpcHits) { return std::min(tpcMaxTrackletHits, hitOffset + tpcHits * tpcTrackletHitsPerHit) * factor; } - size_t NTPCSectorTracks(size_t tpcHits) { return std::min(tpcMaxSectorTracks, offset + tpcHits * tpcSectorTracksPerHit) * factor; } - size_t NTPCSectorTrackHits(size_t tpcHits) { return std::min(tpcMaxSectorTrackHits, offset + tpcHits * tpcSectorTrackHitsPerHit) * factor; } - size_t NTPCMergedTracks(size_t tpcSliceTracks) { return std::min(tpcMaxMergedTracks, offset + tpcSliceTracks * tpcMergedTrackPerSliceTrack) * factor; } - size_t NTPCMergedTrackHits(size_t tpcSliceTrackHitss) { return std::min(tpcMaxMergedTrackHits, offset + tpcSliceTrackHitss * tpcMergedTrackHitPerSliceHit) * factor; } + void rescaleMaxMem(size_t newAvailableMemory); + inline size_t getValue(size_t maxVal, size_t val) + { + return returnMaxVal ? maxVal : (std::min(maxVal, offset + val) * factor); + } + + inline size_t NTPCPeaks(size_t tpcDigits) { return getValue(tpcMaxPeaks, hitOffset + tpcDigits * tpcPeaksPerDigit); } + inline size_t NTPCClusters(size_t tpcDigits) { return getValue(tpcMaxClusters, tpcClustersPerPeak * NTPCPeaks(tpcDigits)); } + inline size_t NTPCStartHits(size_t tpcHits) { return getValue(tpcMaxStartHits, tpcHits * tpcStartHitsPerHit); } + inline size_t NTPCRowStartHits(size_t tpcHits) { return getValue(tpcMaxRowStartHits, NTPCStartHits(tpcHits) / GPUCA_ROW_COUNT * 4.); } + inline size_t NTPCTracklets(size_t tpcHits) { return getValue(tpcMaxTracklets, NTPCStartHits(tpcHits) * tpcTrackletsPerStartHit); } + inline size_t NTPCTrackletHits(size_t tpcHits) { return getValue(tpcMaxTrackletHits, hitOffset + tpcHits * tpcTrackletHitsPerHit); } + inline size_t NTPCSectorTracks(size_t tpcHits) { return getValue(tpcMaxSectorTracks, tpcHits * tpcSectorTracksPerHit); } + inline size_t NTPCSectorTrackHits(size_t tpcHits) { return getValue(tpcMaxSectorTrackHits, tpcHits * tpcSectorTrackHitsPerHit); } + inline size_t NTPCMergedTracks(size_t tpcSliceTracks) { return getValue(tpcMaxMergedTracks, tpcSliceTracks * tpcMergedTrackPerSliceTrack); } + inline size_t NTPCMergedTrackHits(size_t tpcSliceTrackHitss) { return getValue(tpcMaxMergedTrackHits, tpcSliceTrackHitss * tpcMergedTrackHitPerSliceHit); } }; } // namespace GPUCA_NAMESPACE::gpu diff --git a/GPU/GPUTracking/Definitions/GPUSettingsList.h b/GPU/GPUTracking/Definitions/GPUSettingsList.h index e758b76c0dab3..38b16c2a20699 100644 --- a/GPU/GPUTracking/Definitions/GPUSettingsList.h +++ b/GPU/GPUTracking/Definitions/GPUSettingsList.h @@ -117,6 +117,7 @@ AddOption(memoryAllocationStrategy, char, 0, "", 0, "Memory Allocation Stragegy AddOption(forceMemoryPoolSize, unsigned long, 1, "memSize", 0, "Force size of allocated GPU / page locked host memory", min(0ul)) AddOption(forceHostMemoryPoolSize, unsigned long, 0, "hostMemSize", 0, "Force size of allocated host page locked host memory (overriding memSize)", min(0ul)) AddOption(memoryScalingFactor, float, 1.f, "", 0, "Factor to apply to all memory scalers") +AddOption(forceMaxMemScalers, unsigned long, 0, "", 0, "Force using the maximum values for all buffers, Set a value n > 1 to rescale all maximums to a memory size of n") AddOption(registerStandaloneInputMemory, bool, false, "registerInputMemory", 0, "Automatically register input memory buffers for the GPU") AddOption(ompThreads, int, -1, "omp", 't', "Number of OMP threads to run (-1: all)", min(-1), message("Using %s OMP threads")) AddOption(ompKernels, unsigned char, 2, "", 0, "Parallelize with OMP inside kernels instead of over slices, 2 for nested parallelization over TPC sectors and inside kernels") From 8fc95ce7cb6a5f548392c77d3876e9791145ca4c Mon Sep 17 00:00:00 2001 From: David Rohr Date: Tue, 2 Mar 2021 14:39:59 +0100 Subject: [PATCH 04/20] GPU: Fix debug feature checkKernelFailures --- GPU/GPUTracking/Base/GPUReconstructionCPU.cxx | 2 +- GPU/GPUTracking/Base/GPUReconstructionCPU.h | 2 +- GPU/GPUTracking/Base/GPUReconstructionDeviceBase.h | 2 +- GPU/GPUTracking/Base/cuda/GPUReconstructionCUDA.cu | 10 ++++++---- GPU/GPUTracking/Base/cuda/GPUReconstructionCUDA.h | 2 +- GPU/GPUTracking/Base/hip/GPUReconstructionHIP.h | 2 +- GPU/GPUTracking/Base/hip/GPUReconstructionHIP.hip.cxx | 4 ++-- .../Base/opencl-common/GPUReconstructionOCL.cxx | 4 ++-- .../Base/opencl-common/GPUReconstructionOCL.h | 2 +- GPU/GPUTracking/Definitions/GPUSettingsList.h | 2 +- 10 files changed, 17 insertions(+), 15 deletions(-) diff --git a/GPU/GPUTracking/Base/GPUReconstructionCPU.cxx b/GPU/GPUTracking/Base/GPUReconstructionCPU.cxx index a0991b838dfea..7c84fcda63227 100644 --- a/GPU/GPUTracking/Base/GPUReconstructionCPU.cxx +++ b/GPU/GPUTracking/Base/GPUReconstructionCPU.cxx @@ -112,7 +112,7 @@ size_t GPUReconstructionCPU::GPUMemCpyAlways(bool onGpu, void* dst, const void* return 0; } size_t GPUReconstructionCPU::WriteToConstantMemory(size_t offset, const void* src, size_t size, int stream, deviceEvent* ev) { return 0; } -int GPUReconstructionCPU::GPUDebug(const char* state, int stream) { return 0; } +int GPUReconstructionCPU::GPUDebug(const char* state, int stream, bool force) { return 0; } size_t GPUReconstructionCPU::TransferMemoryResourcesHelper(GPUProcessor* proc, int stream, bool all, bool toGPU) { int inc = toGPU ? GPUMemoryResource::MEMORY_INPUT_FLAG : GPUMemoryResource::MEMORY_OUTPUT_FLAG; diff --git a/GPU/GPUTracking/Base/GPUReconstructionCPU.h b/GPU/GPUTracking/Base/GPUReconstructionCPU.h index dfc26366e2793..29eb203e7713c 100644 --- a/GPU/GPUTracking/Base/GPUReconstructionCPU.h +++ b/GPU/GPUTracking/Base/GPUReconstructionCPU.h @@ -137,7 +137,7 @@ class GPUReconstructionCPU : public GPUReconstructionKernels constexpr static const char* GetKernelName(); - virtual int GPUDebug(const char* state = "UNKNOWN", int stream = -1); + virtual int GPUDebug(const char* state = "UNKNOWN", int stream = -1, bool force = false); int registerMemoryForGPU(const void* ptr, size_t size) override { return 0; } int unregisterMemoryForGPU(const void* ptr) override { return 0; } int GPUStuck() { return mGPUStuck; } diff --git a/GPU/GPUTracking/Base/GPUReconstructionDeviceBase.h b/GPU/GPUTracking/Base/GPUReconstructionDeviceBase.h index 7e81e7b59151a..12080b895cdf6 100644 --- a/GPU/GPUTracking/Base/GPUReconstructionDeviceBase.h +++ b/GPU/GPUTracking/Base/GPUReconstructionDeviceBase.h @@ -46,7 +46,7 @@ class GPUReconstructionDeviceBase : public GPUReconstructionCPU virtual const GPUTPCTracker* CPUTracker(int iSlice) { return &processors()->tpcTrackers[iSlice]; } - int GPUDebug(const char* state = "UNKNOWN", int stream = -1) override = 0; + int GPUDebug(const char* state = "UNKNOWN", int stream = -1, bool force = false) override = 0; size_t TransferMemoryInternal(GPUMemoryResource* res, int stream, deviceEvent* ev, deviceEvent* evList, int nEvents, bool toGPU, const void* src, void* dst) override = 0; size_t GPUMemCpy(void* dst, const void* src, size_t size, int stream, int toGPU, deviceEvent* ev = nullptr, deviceEvent* evList = nullptr, int nEvents = 1) override = 0; size_t GPUMemCpyAlways(bool onGpu, void* dst, const void* src, size_t size, int stream, int toGPU, deviceEvent* ev = nullptr, deviceEvent* evList = nullptr, int nEvents = 1) override; diff --git a/GPU/GPUTracking/Base/cuda/GPUReconstructionCUDA.cu b/GPU/GPUTracking/Base/cuda/GPUReconstructionCUDA.cu index 5cdfdbdf361a1..aea0199646d70 100644 --- a/GPU/GPUTracking/Base/cuda/GPUReconstructionCUDA.cu +++ b/GPU/GPUTracking/Base/cuda/GPUReconstructionCUDA.cu @@ -171,7 +171,9 @@ void GPUReconstructionCUDABackend::runKernelBackendInternal(krnlSetup& _xyz, con backendInternal::runKernelBackendMacro(_xyz, this, args...); } if (mProcessingSettings.checkKernelFailures) { - GPUDebug(GetKernelName(), _xyz.x.stream); + if (GPUDebug(GetKernelName(), _xyz.x.stream, true)) { + throw std::runtime_error("Kernel Failure"); + } } } @@ -705,7 +707,7 @@ bool GPUReconstructionCUDABackend::IsEventDone(deviceEvent* evList, int nEvents) return (true); } -int GPUReconstructionCUDABackend::GPUDebug(const char* state, int stream) +int GPUReconstructionCUDABackend::GPUDebug(const char* state, int stream, bool force) { // Wait for CUDA-Kernel to finish and check for CUDA errors afterwards, in case of debugmode cudaError cuErr; @@ -714,10 +716,10 @@ int GPUReconstructionCUDABackend::GPUDebug(const char* state, int stream) GPUError("Cuda Error %s while running kernel (%s) (Stream %d)", cudaGetErrorString(cuErr), state, stream); return (1); } - if (mProcessingSettings.debugLevel <= 0) { + if (force == false && mProcessingSettings.debugLevel <= 0) { return (0); } - if (GPUFailedMsgI(cudaDeviceSynchronize())) { + if (GPUFailedMsgI(stream == -1 ? cudaDeviceSynchronize() : cudaStreamSynchronize(mInternals->Streams[stream]))) { GPUError("CUDA Error while synchronizing (%s) (Stream %d)", state, stream); return (1); } diff --git a/GPU/GPUTracking/Base/cuda/GPUReconstructionCUDA.h b/GPU/GPUTracking/Base/cuda/GPUReconstructionCUDA.h index d3188865df4b6..69982067b7686 100644 --- a/GPU/GPUTracking/Base/cuda/GPUReconstructionCUDA.h +++ b/GPU/GPUTracking/Base/cuda/GPUReconstructionCUDA.h @@ -53,7 +53,7 @@ class GPUReconstructionCUDABackend : public GPUReconstructionDeviceBase std::unique_ptr GetThreadContext() override; bool CanQueryMaxMemory() override { return true; } void SynchronizeGPU() override; - int GPUDebug(const char* state = "UNKNOWN", int stream = -1) override; + int GPUDebug(const char* state = "UNKNOWN", int stream = -1, bool force = false) override; void SynchronizeStream(int stream) override; void SynchronizeEvents(deviceEvent* evList, int nEvents = 1) override; void StreamWaitForEvents(int stream, deviceEvent* evList, int nEvents = 1) override; diff --git a/GPU/GPUTracking/Base/hip/GPUReconstructionHIP.h b/GPU/GPUTracking/Base/hip/GPUReconstructionHIP.h index 9f36d0db0fb78..d24f5129c2af6 100644 --- a/GPU/GPUTracking/Base/hip/GPUReconstructionHIP.h +++ b/GPU/GPUTracking/Base/hip/GPUReconstructionHIP.h @@ -41,7 +41,7 @@ class GPUReconstructionHIPBackend : public GPUReconstructionDeviceBase void UpdateSettings() override; void SynchronizeGPU() override; - int GPUDebug(const char* state = "UNKNOWN", int stream = -1) override; + int GPUDebug(const char* state = "UNKNOWN", int stream = -1, bool force = false) override; void SynchronizeStream(int stream) override; void SynchronizeEvents(deviceEvent* evList, int nEvents = 1) override; void StreamWaitForEvents(int stream, deviceEvent* evList, int nEvents = 1) override; diff --git a/GPU/GPUTracking/Base/hip/GPUReconstructionHIP.hip.cxx b/GPU/GPUTracking/Base/hip/GPUReconstructionHIP.hip.cxx index 432c6ccbf4419..f6ea49e51f311 100644 --- a/GPU/GPUTracking/Base/hip/GPUReconstructionHIP.hip.cxx +++ b/GPU/GPUTracking/Base/hip/GPUReconstructionHIP.hip.cxx @@ -542,7 +542,7 @@ bool GPUReconstructionHIPBackend::IsEventDone(deviceEvent* evList, int nEvents) return (true); } -int GPUReconstructionHIPBackend::GPUDebug(const char* state, int stream) +int GPUReconstructionHIPBackend::GPUDebug(const char* state, int stream, bool force) { // Wait for HIP-Kernel to finish and check for HIP errors afterwards, in case of debugmode hipError_t cuErr; @@ -551,7 +551,7 @@ int GPUReconstructionHIPBackend::GPUDebug(const char* state, int stream) GPUError("HIP Error %s while running kernel (%s) (Stream %d)", hipGetErrorString(cuErr), state, stream); return (1); } - if (mProcessingSettings.debugLevel <= 0) { + if (!force && mProcessingSettings.debugLevel <= 0) { return (0); } if (GPUFailedMsgI(hipDeviceSynchronize())) { diff --git a/GPU/GPUTracking/Base/opencl-common/GPUReconstructionOCL.cxx b/GPU/GPUTracking/Base/opencl-common/GPUReconstructionOCL.cxx index 69fd3d5bba07e..aff6dc88c4703 100644 --- a/GPU/GPUTracking/Base/opencl-common/GPUReconstructionOCL.cxx +++ b/GPU/GPUTracking/Base/opencl-common/GPUReconstructionOCL.cxx @@ -462,10 +462,10 @@ bool GPUReconstructionOCL::IsEventDone(deviceEvent* evList, int nEvents) return true; } -int GPUReconstructionOCL::GPUDebug(const char* state, int stream) +int GPUReconstructionOCL::GPUDebug(const char* state, int stream, bool force) { // Wait for OPENCL-Kernel to finish and check for OPENCL errors afterwards, in case of debugmode - if (mProcessingSettings.debugLevel <= 0) { + if (!force && mProcessingSettings.debugLevel <= 0) { return (0); } for (int i = 0; i < mNStreams; i++) { diff --git a/GPU/GPUTracking/Base/opencl-common/GPUReconstructionOCL.h b/GPU/GPUTracking/Base/opencl-common/GPUReconstructionOCL.h index 46e50e8547357..3ada2b683aa69 100644 --- a/GPU/GPUTracking/Base/opencl-common/GPUReconstructionOCL.h +++ b/GPU/GPUTracking/Base/opencl-common/GPUReconstructionOCL.h @@ -39,7 +39,7 @@ class GPUReconstructionOCL : public GPUReconstructionDeviceBase void SynchronizeGPU() override; int DoStuckProtection(int stream, void* event) override; - int GPUDebug(const char* state = "UNKNOWN", int stream = -1) override; + int GPUDebug(const char* state = "UNKNOWN", int stream = -1, bool force = false) override; void SynchronizeStream(int stream) override; void SynchronizeEvents(deviceEvent* evList, int nEvents = 1) override; void StreamWaitForEvents(int stream, deviceEvent* evList, int nEvents = 1) override; diff --git a/GPU/GPUTracking/Definitions/GPUSettingsList.h b/GPU/GPUTracking/Definitions/GPUSettingsList.h index 38b16c2a20699..e6346aff5d70c 100644 --- a/GPU/GPUTracking/Definitions/GPUSettingsList.h +++ b/GPU/GPUTracking/Definitions/GPUSettingsList.h @@ -104,7 +104,7 @@ AddOption(trdNMaxCollisions, int, 1000, "", 0, "Maximum number of collisions per AddOption(debugLevel, int, -1, "debug", 'd', "Set debug level (-1 = silend)") AddOption(allocDebugLevel, int, 0, "allocDebug", 0, "Some debug output for memory allocations (without messing with normal debug level)") AddOption(debugMask, int, -1, "", 0, "Mask for debug output dumps to file") -AddOption(checkKernelFailures, int, 0, "", 0, "Synchronize after each kernel call and identify failing kernels") +AddOption(checkKernelFailures, bool, false, "", 0, "Synchronize after each kernel call and identify failing kernels") AddOption(comparableDebutOutput, bool, true, "", 0, "Make CPU and GPU debug output comparable (sort / skip concurrent parts)") AddOption(showOutputStat, bool, false, "", 0, "Print some track output statistics") AddOption(runCompressionStatistics, bool, false, "compressionStat", 0, "Run statistics and verification for cluster compression") From 789558aaa8a633bb11ffae4fdea659bfe6784dad Mon Sep 17 00:00:00 2001 From: David Rohr Date: Tue, 2 Mar 2021 15:05:34 +0100 Subject: [PATCH 05/20] GPU: Add error handling in case merged track buffer overflows --- GPU/GPUTracking/Global/GPUErrorCodes.h | 2 ++ GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx | 14 ++++++++++++-- 2 files changed, 14 insertions(+), 2 deletions(-) diff --git a/GPU/GPUTracking/Global/GPUErrorCodes.h b/GPU/GPUTracking/Global/GPUErrorCodes.h index b7b6df0938d88..07a616927e0ef 100644 --- a/GPU/GPUTracking/Global/GPUErrorCodes.h +++ b/GPU/GPUTracking/Global/GPUErrorCodes.h @@ -28,3 +28,5 @@ GPUCA_ERROR_CODE(12, ERROR_SLICEDATA_FIRSTHITINBIN_OVERFLOW) GPUCA_ERROR_CODE(13, ERROR_SLICEDATA_HITINROW_OVERFLOW) GPUCA_ERROR_CODE(14, ERROR_SLICEDATA_BIN_OVERFLOW) GPUCA_ERROR_CODE(15, ERROR_SLICEDATA_Z_OVERFLOW) +GPUCA_ERROR_CODE(16, ERROR_MERGER_HIT_OVERFLOW) +GPUCA_ERROR_CODE(17, ERROR_MERGER_TRACK_OVERFLOW) diff --git a/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx b/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx index 5d137bb8a04a3..cc4381aa7a139 100644 --- a/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx +++ b/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx @@ -1664,7 +1664,12 @@ GPUd() void GPUTPCGMMerger::CollectMergedTracks(int nBlocks, int nThreads, int i nHits = nFilteredHits; } - int iOutTrackFirstCluster = CAMath::AtomicAdd(&mMemory->nOutputTrackClusters, (unsigned int)nHits); + unsigned int iOutTrackFirstCluster = CAMath::AtomicAdd(&mMemory->nOutputTrackClusters, (unsigned int)nHits); + if (iOutTrackFirstCluster >= mNMaxOutputTrackClusters) { + raiseError(GPUErrors::ERROR_MERGER_HIT_OVERFLOW, iOutTrackFirstCluster, mNMaxOutputTrackClusters); + CAMath::AtomicExch(&mMemory->nOutputTrackClusters, 0u); + continue; + } GPUTPCGMMergedTrackHit* cl = mClusters + iOutTrackFirstCluster; GPUTPCGMMergedTrackHitXYZ* clXYZ = mClustersXYZ + iOutTrackFirstCluster; @@ -1717,7 +1722,12 @@ GPUd() void GPUTPCGMMerger::CollectMergedTracks(int nBlocks, int nThreads, int i cl[i].leg = trackClusters[i].leg; } // nHits - int iOutputTrack = CAMath::AtomicAdd(&mMemory->nOutputTracks, 1u); + unsigned int iOutputTrack = CAMath::AtomicAdd(&mMemory->nOutputTracks, 1u); + if (iOutputTrack >= mNMaxTracks) { + raiseError(GPUErrors::ERROR_MERGER_TRACK_OVERFLOW, iOutputTrack, mNMaxTracks); + CAMath::AtomicExch(&mMemory->nOutputTracks, 0u); + continue; + } GPUTPCGMMergedTrack& mergedTrack = mOutputTracks[iOutputTrack]; From a727d1f73b3330bce6e95eabcb72442f849def4c Mon Sep 17 00:00:00 2001 From: David Rohr Date: Tue, 2 Mar 2021 15:51:41 +0100 Subject: [PATCH 06/20] GPU: Improve debug output on memory usage --- GPU/GPUTracking/Base/GPUReconstruction.cxx | 10 +++++++--- .../Global/GPUChainTrackingDebugAndProfiling.cxx | 15 ++++++++++----- 2 files changed, 17 insertions(+), 8 deletions(-) diff --git a/GPU/GPUTracking/Base/GPUReconstruction.cxx b/GPU/GPUTracking/Base/GPUReconstruction.cxx index 7bfb66b55f419..1db577b8b31f2 100644 --- a/GPU/GPUTracking/Base/GPUReconstruction.cxx +++ b/GPU/GPUTracking/Base/GPUReconstruction.cxx @@ -623,7 +623,7 @@ size_t GPUReconstruction::AllocateRegisteredMemory(short ires, GPUOutputControl* void* GPUReconstruction::AllocateUnmanagedMemory(size_t size, int type) { if (type != GPUMemoryResource::MEMORY_HOST && (!IsGPU() || type != GPUMemoryResource::MEMORY_GPU)) { - throw std::bad_alloc(); + throw std::runtime_error("Requested invalid memory typo for unmanaged allocation"); } if (mProcessingSettings.memoryAllocationStrategy == GPUMemoryResource::ALLOCATION_INDIVIDUAL) { mUnmanagedChunks.emplace_back(new char[size + GPUCA_BUFFER_ALIGNMENT]); @@ -634,6 +634,7 @@ void* GPUReconstruction::AllocateUnmanagedMemory(size_t size, int type) char* retVal; GPUProcessor::computePointerWithAlignment(pool, retVal, size); if (pool > poolend) { + GPUError("Insufficient unmanaged memory: missing %lu", (size_t)((char*)pool - (char*)poolend)); throw std::bad_alloc(); } UpdateMaxMemoryUsed(); @@ -652,6 +653,7 @@ void* GPUReconstruction::AllocateVolatileDeviceMemory(size_t size) char* retVal; GPUProcessor::computePointerWithAlignment(mDeviceMemoryPool, retVal, size); if (mDeviceMemoryPool > mDeviceMemoryPoolEnd) { + GPUError("Insufficient volatile device memory: missing %lu", (size_t)((char*)mDeviceMemoryPool - (char*)mDeviceMemoryPoolEnd)); throw std::bad_alloc(); } UpdateMaxMemoryUsed(); @@ -731,7 +733,9 @@ void GPUReconstruction::PopNonPersistentMemory(RecoStep step) if (IsGPU()) { printf("Allocated Device memory after %30s: %'13lld (non temporary %'13lld, blocked %'13lld)\n", GPUDataTypes::RECO_STEP_NAMES[getRecoStepNum(step, true)], ptrDiff(mDeviceMemoryPool, mDeviceMemoryBase) + ptrDiff((char*)mDeviceMemoryBase + mDeviceMemorySize, mDeviceMemoryPoolEnd), ptrDiff(mDeviceMemoryPool, mDeviceMemoryBase), mDeviceMemoryPoolBlocked == nullptr ? 0ll : ptrDiff((char*)mDeviceMemoryBase + mDeviceMemorySize, mDeviceMemoryPoolBlocked)); } - printf("Allocated Host memory after %30s: %'13lld (non temporary %'13lld, blocked %'13lld)\n", GPUDataTypes::RECO_STEP_NAMES[getRecoStepNum(step, true)], ptrDiff(mHostMemoryPool, mHostMemoryBase) + ptrDiff((char*)mHostMemoryBase + mHostMemorySize, mHostMemoryPoolEnd), ptrDiff(mHostMemoryPool, mHostMemoryBase), mHostMemoryPoolBlocked == nullptr ? 0ll : ptrDiff((char*)mHostMemoryBase + mHostMemorySize, mHostMemoryPoolBlocked)); + printf("Allocated Host memory after %30s: %'13lld (non temporary %'13lld, blocked %'13lld)\n", GPUDataTypes::RECO_STEP_NAMES[getRecoStepNum(step, true)], ptrDiff(mHostMemoryPool, mHostMemoryBase) + ptrDiff((char*)mHostMemoryBase + mHostMemorySize, mHostMemoryPoolEnd), ptrDiff(mHostMemoryPool, mHostMemoryBase), mHostMemoryPoolBlocked == nullptr ? 0ll : ptrDiff((char*)mHostMemoryBase + mHostMemorySize, mHostMemoryPoolBlocked)); + printf("%16s", ""); + PrintMemoryMax(); } if (mProcessingSettings.keepDisplayMemory || mProcessingSettings.disableMemoryReuse) { return; @@ -833,7 +837,7 @@ void GPUReconstruction::PrintMemoryStatistics() } printf("%59s CPU / %9s GPU\n", "", ""); for (auto it = sizes.begin(); it != sizes.end(); it++) { - printf("Allocation %30s %s: Size %'13lld / %'13lld\n", it->first.c_str(), it->second[2] ? "P" : " ", (long long int)it->second[0], (long long int)it->second[1]); + printf("Allocation %30s %s: Size %'14lld / %'14lld\n", it->first.c_str(), it->second[2] ? "P" : " ", (long long int)it->second[0], (long long int)it->second[1]); } PrintMemoryOverview(); for (unsigned int i = 0; i < mChains.size(); i++) { diff --git a/GPU/GPUTracking/Global/GPUChainTrackingDebugAndProfiling.cxx b/GPU/GPUTracking/Global/GPUChainTrackingDebugAndProfiling.cxx index 2a04653a68f85..1083f34b7ebed 100644 --- a/GPU/GPUTracking/Global/GPUChainTrackingDebugAndProfiling.cxx +++ b/GPU/GPUTracking/Global/GPUChainTrackingDebugAndProfiling.cxx @@ -13,6 +13,7 @@ #include "GPUChainTracking.h" #include "GPUTrackingInputProvider.h" +#include "GPUMemorySizeScalers.h" #include #include #include @@ -139,18 +140,22 @@ void GPUChainTracking::PrintMemoryStatistics() addToMap("TPC Clusterer Sector Peaks", usageMap, processors()->tpcClusterer[i].mPmemory->counters.nPeaks, processors()->tpcClusterer[i].mNMaxPeaks); addToMap("TPC Clusterer Sector Clusters", usageMap, processors()->tpcClusterer[i].mPmemory->counters.nClusters, processors()->tpcClusterer[i].mNMaxClusters); #endif - addToMap("TPC Start Hits", usageMap, *processors()->tpcTrackers[i].NStartHits(), processors()->tpcTrackers[i].NMaxStartHits()); - addToMap("TPC Tracklets", usageMap, *processors()->tpcTrackers[i].NTracklets(), processors()->tpcTrackers[i].NMaxTracklets()); - addToMap("TPC TrackletHits", usageMap, *processors()->tpcTrackers[i].NRowHits(), processors()->tpcTrackers[i].NMaxRowHits()); + addToMap("TPC Sector Start Hits", usageMap, *processors()->tpcTrackers[i].NStartHits(), processors()->tpcTrackers[i].NMaxStartHits()); + addToMap("TPC Sector Tracklets", usageMap, *processors()->tpcTrackers[i].NTracklets(), processors()->tpcTrackers[i].NMaxTracklets()); + addToMap("TPC Sector TrackletHits", usageMap, *processors()->tpcTrackers[i].NRowHits(), processors()->tpcTrackers[i].NMaxRowHits()); addToMap("TPC Sector Tracks", usageMap, *processors()->tpcTrackers[i].NTracks(), processors()->tpcTrackers[i].NMaxTracks()); addToMap("TPC Sector TrackHits", usageMap, *processors()->tpcTrackers[i].NTrackHits(), processors()->tpcTrackers[i].NMaxTrackHits()); } - addToMap("TPC Clusters", usageMap, mIOPtrs.clustersNative->nClustersTotal, mInputsHost->mNClusterNative); + addToMap("TPC Clusterer Clusters", usageMap, mRec->MemoryScalers()->nTPCHits, mRec->MemoryScalers()->NTPCClusters(mRec->MemoryScalers()->nTPCdigits)); addToMap("TPC Tracks", usageMap, processors()->tpcMerger.NOutputTracks(), processors()->tpcMerger.NMaxTracks()); addToMap("TPC TrackHits", usageMap, processors()->tpcMerger.NOutputTrackClusters(), processors()->tpcMerger.NMaxOutputTrackClusters()); + if (mRec->GetProcessingSettings().createO2Output) { + addToMap("TPC O2 Tracks", usageMap, processors()->tpcMerger.NOutputTracksTPCO2(), processors()->tpcMerger.NOutputTracksTPCO2()); + addToMap("TPC O2 ClusRefs", usageMap, processors()->tpcMerger.NOutputClusRefsTPCO2(), processors()->tpcMerger.NOutputClusRefsTPCO2()); + } for (auto& elem : usageMap) { - GPUInfo("Mem Usage %-30s : %9lu / %9lu (%3.0f%% / %3.0f%% / count %3u / max %9lu)", elem.first.c_str(), elem.second.nSum, elem.second.nBoundSum, 100. * elem.second.nSum / std::max(1lu, elem.second.nBoundSum), 100. * elem.second.maxUse, elem.second.count, elem.second.nMax); + printf("Mem Usage %-30s : %'14lu / %'14lu (%3.0f%% / %3.0f%% / count %3u / max %'14lu)\n", elem.first.c_str(), elem.second.nSum, elem.second.nBoundSum, 100. * elem.second.nSum / std::max(1lu, elem.second.nBoundSum), 100. * elem.second.maxUse, elem.second.count, elem.second.nMax); } } From bd0bb619ebdf5b275456cf388ab7add284a9b9ec Mon Sep 17 00:00:00 2001 From: David Rohr Date: Tue, 2 Mar 2021 19:39:11 +0100 Subject: [PATCH 07/20] GPU: Tune GPU-memory depending memory cuts --- .../DataTypes/GPUMemorySizeScalers.cxx | 21 ++++++++++--------- .../DataTypes/GPUMemorySizeScalers.h | 18 ++++++++-------- 2 files changed, 20 insertions(+), 19 deletions(-) diff --git a/GPU/GPUTracking/DataTypes/GPUMemorySizeScalers.cxx b/GPU/GPUTracking/DataTypes/GPUMemorySizeScalers.cxx index 95edfdba17841..4f73b064e0783 100644 --- a/GPU/GPUTracking/DataTypes/GPUMemorySizeScalers.cxx +++ b/GPU/GPUTracking/DataTypes/GPUMemorySizeScalers.cxx @@ -18,15 +18,16 @@ using namespace GPUCA_NAMESPACE::gpu; void GPUMemorySizeScalers::rescaleMaxMem(size_t newAvailableMemory) { GPUMemorySizeScalers tmp; - tpcMaxPeaks = (double)tmp.tpcMaxPeaks * newAvailableMemory / tmp.availableMemory; - tpcMaxClusters = (double)tmp.tpcMaxClusters * newAvailableMemory / tmp.availableMemory; - tpcMaxStartHits = (double)tmp.tpcMaxStartHits * newAvailableMemory / tmp.availableMemory; - tpcMaxRowStartHits = (double)tmp.tpcMaxRowStartHits * newAvailableMemory / tmp.availableMemory; - tpcMaxTracklets = (double)tmp.tpcMaxTracklets * newAvailableMemory / tmp.availableMemory; - tpcMaxTrackletHits = (double)tmp.tpcMaxTrackletHits * newAvailableMemory / tmp.availableMemory; - tpcMaxSectorTracks = (double)tmp.tpcMaxSectorTracks * newAvailableMemory / tmp.availableMemory; - tpcMaxSectorTrackHits = (double)tmp.tpcMaxSectorTrackHits * newAvailableMemory / tmp.availableMemory; - tpcMaxMergedTracks = (double)tmp.tpcMaxMergedTracks * newAvailableMemory / tmp.availableMemory; - tpcMaxMergedTrackHits = (double)tmp.tpcMaxMergedTrackHits * newAvailableMemory / tmp.availableMemory; + double scaleFactor = newAvailableMemory > tmp.availableMemory ? ((double)newAvailableMemory / tmp.availableMemory) : 1.; + tpcMaxPeaks = (double)tmp.tpcMaxPeaks * scaleFactor; + tpcMaxClusters = (double)tmp.tpcMaxClusters * scaleFactor; + tpcMaxStartHits = (double)tmp.tpcMaxStartHits * scaleFactor; + tpcMaxRowStartHits = (double)tmp.tpcMaxRowStartHits * scaleFactor; + tpcMaxTracklets = (double)tmp.tpcMaxTracklets * scaleFactor; + tpcMaxTrackletHits = (double)tmp.tpcMaxTrackletHits * scaleFactor; + tpcMaxSectorTracks = (double)tmp.tpcMaxSectorTracks * scaleFactor; + tpcMaxSectorTrackHits = (double)tmp.tpcMaxSectorTrackHits * scaleFactor; + tpcMaxMergedTracks = (double)tmp.tpcMaxMergedTracks * scaleFactor; + tpcMaxMergedTrackHits = (double)tmp.tpcMaxMergedTrackHits * scaleFactor; availableMemory = newAvailableMemory; } diff --git a/GPU/GPUTracking/DataTypes/GPUMemorySizeScalers.h b/GPU/GPUTracking/DataTypes/GPUMemorySizeScalers.h index 7e1d3e786ceae..ed0aab483943b 100644 --- a/GPU/GPUTracking/DataTypes/GPUMemorySizeScalers.h +++ b/GPU/GPUTracking/DataTypes/GPUMemorySizeScalers.h @@ -46,16 +46,16 @@ struct GPUMemorySizeScalers { // Upper limits size_t tpcMaxPeaks = 1000000000; - size_t tpcMaxClusters = 620000000; - size_t tpcMaxStartHits = 1250000; + size_t tpcMaxClusters = 320000000; + size_t tpcMaxStartHits = 650000; size_t tpcMaxRowStartHits = 1000000000; - size_t tpcMaxTracklets = 1000000; - size_t tpcMaxTrackletHits = 66000000; - size_t tpcMaxSectorTracks = 250000; - size_t tpcMaxSectorTrackHits = 11500000; - size_t tpcMaxMergedTracks = 5800000; - size_t tpcMaxMergedTrackHits = 380000000; - size_t availableMemory = 24000000000; + size_t tpcMaxTracklets = 520000; + size_t tpcMaxTrackletHits = 35000000; + size_t tpcMaxSectorTracks = 130000; + size_t tpcMaxSectorTrackHits = 5900000; + size_t tpcMaxMergedTracks = 3000000; + size_t tpcMaxMergedTrackHits = 200000000; + size_t availableMemory = 22000000000; bool returnMaxVal = false; void rescaleMaxMem(size_t newAvailableMemory); From c39627d47b5efe5bb5125ec8dc7da5419e8db493 Mon Sep 17 00:00:00 2001 From: David Rohr Date: Tue, 2 Mar 2021 19:40:03 +0100 Subject: [PATCH 08/20] GPU: Clean up some manual pointer computation in TPC track merging --- .../Global/GPUChainTrackingMerger.cxx | 3 +- GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx | 68 +++++++++---------- GPU/GPUTracking/Merger/GPUTPCGMMerger.h | 16 ++++- GPU/GPUTracking/Merger/GPUTPCGMO2Output.cxx | 19 ++---- 4 files changed, 56 insertions(+), 50 deletions(-) diff --git a/GPU/GPUTracking/Global/GPUChainTrackingMerger.cxx b/GPU/GPUTracking/Global/GPUChainTrackingMerger.cxx index cdaceef4d5d00..55ec254b4668c 100644 --- a/GPU/GPUTracking/Global/GPUChainTrackingMerger.cxx +++ b/GPU/GPUTracking/Global/GPUChainTrackingMerger.cxx @@ -185,8 +185,7 @@ int GPUChainTracking::RunTPCTrackingMerger(bool synchronizeOutput) throw std::runtime_error("mNMaxClusters too small"); } if (!param().rec.NonConsecutiveIDs) { - unsigned int* sharedCount = (unsigned int*)MergerShadowAll.TmpMem() + CAMath::nextMultipleOf<4>(Merger.Memory()->nOutputTracks); - runKernel({numBlocks, -ThreadCount(), 0, deviceType, RecoStep::TPCMerging}, krnlRunRangeNone, {}, sharedCount, maxId * sizeof(*sharedCount)); + runKernel({numBlocks, -ThreadCount(), 0, deviceType, RecoStep::TPCMerging}, krnlRunRangeNone, {}, MergerShadowAll.SharedCount(), maxId * sizeof(*MergerShadowAll.SharedCount())); runKernel({numBlocks, -ThreadCount(), 0, deviceType, RecoStep::TPCMerging}, krnlRunRangeNone, {}, MergerShadowAll.ClusterAttachment(), maxId * sizeof(*MergerShadowAll.ClusterAttachment())); runKernel(GetGridAuto(0, deviceType), krnlRunRangeNone, krnlEventNone); CondWaitEvent(waitForTransfer, &mEvents->single); diff --git a/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx b/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx index cc4381aa7a139..da6cee561adb0 100644 --- a/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx +++ b/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx @@ -75,7 +75,7 @@ static constexpr int kMaxClusters = GPUCA_MERGER_MAX_TRACK_CLUSTERS; #include "GPUMemorySizeScalers.h" GPUTPCGMMerger::GPUTPCGMMerger() - : mTrackLinks(nullptr), mNMaxSliceTracks(0), mNMaxTracks(0), mNMaxSingleSliceTracks(0), mNMaxOutputTrackClusters(0), mNMaxClusters(0), mMemoryResMemory(-1), mNClusters(0), mOutputTracks(nullptr), mSliceTrackInfos(nullptr), mSliceTrackInfoIndex(nullptr), mClusters(nullptr), mClustersXYZ(nullptr), mGlobalClusterIDs(nullptr), mClusterAttachment(nullptr), mOutputTracksTPCO2(nullptr), mOutputClusRefsTPCO2(nullptr), mOutputTracksTPCO2MC(nullptr), mTrackOrderAttach(nullptr), mTrackOrderProcess(nullptr), mTmpMem(nullptr), mBorderMemory(nullptr), mBorderRangeMemory(nullptr), mMemory(nullptr), mRetryRefitIds(nullptr), mLoopData(nullptr) + : mTrackLinks(nullptr), mNMaxSliceTracks(0), mNMaxTracks(0), mNMaxSingleSliceTracks(0), mNMaxOutputTrackClusters(0), mNMaxClusters(0), mMemoryResMemory(-1), mNClusters(0), mOutputTracks(nullptr), mSliceTrackInfos(nullptr), mSliceTrackInfoIndex(nullptr), mClusters(nullptr), mClustersXYZ(nullptr), mGlobalClusterIDs(nullptr), mClusterAttachment(nullptr), mOutputTracksTPCO2(nullptr), mOutputClusRefsTPCO2(nullptr), mOutputTracksTPCO2MC(nullptr), mTrackOrderAttach(nullptr), mTrackOrderProcess(nullptr), mBorderMemory(nullptr), mBorderRangeMemory(nullptr), mMemory(nullptr), mRetryRefitIds(nullptr), mLoopData(nullptr) { //* constructor @@ -248,16 +248,9 @@ void* GPUTPCGMMerger::SetPointersMerger(void* mem) if (mRec->GetParam().rec.NonConsecutiveIDs) { computePointerWithAlignment(mem, mGlobalClusterIDs, mNMaxOutputTrackClusters); } + computePointerWithAlignment(mem, mBorderMemory, 2 * mNMaxSliceTracks); computePointerWithAlignment(mem, mBorderRangeMemory, 2 * mNMaxSliceTracks); - computePointerWithAlignment(mem, mTrackLinks, mNMaxSliceTracks); - computePointerWithAlignment(mem, mTrackCCRoots, mNMaxSliceTracks); - size_t tmpSize = CAMath::Max(mNMaxSingleSliceTracks, 1u) * NSLICES * sizeof(int); - tmpSize = CAMath::Max(tmpSize, CAMath::nextMultipleOf<4>(mNMaxTracks) * sizeof(int) + mNMaxClusters * sizeof(unsigned int)); - tmpSize = CAMath::Max(tmpSize, 4 * mNMaxTracks * sizeof(int)); - tmpSize = CAMath::Max(tmpSize, mNMaxTracks * (sizeof(*mRetryRefitIds) + sizeof(*mLoopData))); - computePointerWithAlignment(mem, mTmpMem, (tmpSize + sizeof(*mTmpMem) - 1) / sizeof(*mTmpMem)); - int nTracks = 0; for (int iSlice = 0; iSlice < NSLICES; iSlice++) { const int n = mRec->GetParam().rec.mergerReadFromTrackerDirectly ? *mRec->GetConstantMem().tpcTrackers[iSlice].NTracks() : mkSlices[iSlice]->NTracks(); @@ -266,9 +259,26 @@ void* GPUTPCGMMerger::SetPointersMerger(void* mem) mBorderRange[iSlice] = mBorderRangeMemory + 2 * nTracks; nTracks += n; } - mLoopData = (GPUTPCGMLoopData*)mTmpMem; - mRetryRefitIds = (unsigned int*)(mLoopData + mNMaxTracks); - return mem; + + computePointerWithAlignment(mem, mTrackLinks, mNMaxSliceTracks); + computePointerWithAlignment(mem, mTrackCCRoots, mNMaxSliceTracks); + + void* memBase = mem; + computePointerWithAlignment(mem, mTrackSortO2, mNMaxTracks); // GPUTPCGMO2Output::prepare - GPUTPCGMO2Output::sort - GPUTPCGMO2Output::output + computePointerWithAlignment(mem, mClusRefTmp, mNMaxTracks); + void* memMax = mem; + mem = memBase; + computePointerWithAlignment(mem, mTrackIDs, mNMaxTracks); // UnpackResetIds - RefitSliceTracks - UnpackSliceGlobal + memMax = std::max(mem, memMax); + mem = memBase; + computePointerWithAlignment(mem, mTrackSort, mNMaxTracks); // PrepareClustersForFit0 - SortTracksQPt - PrepareClustersForFit1 - PrepareClustersForFit1 / Finalize0 - Finalize2 + computePointerWithAlignment(mem, mSharedCount, mNMaxClusters); + memMax = std::max(mem, memMax); + mem = memBase; + computePointerWithAlignment(mem, mLoopData, mNMaxTracks); // GPUTPCGMMergerTrackFit - GPUTPCGMMergerFollowLoopers + computePointerWithAlignment(mem, mRetryRefitIds, mNMaxTracks); + memMax = std::max(mem, memMax); + return memMax; } void* GPUTPCGMMerger::SetPointersMemory(void* mem) @@ -511,7 +521,6 @@ GPUd() void GPUTPCGMMerger::UnpackSaveNumber(int id) GPUd() void GPUTPCGMMerger::UnpackSliceGlobal(int nBlocks, int nThreads, int iBlock, int iThread, int iSlice) { - int* TrackIds = (int*)mTmpMem; const GPUTPCTracker& trk = GetConstantMem()->tpcTrackers[iSlice]; float alpha = Param().Alpha(iSlice); const GPUTPCTrack* sliceTr = mMemory->firstGlobalTracks[iSlice]; @@ -523,7 +532,7 @@ GPUd() void GPUTPCGMMerger::UnpackSliceGlobal(int nBlocks, int nThreads, int iBl } else if (itr > nLocalTracks) { sliceTr = sliceTr->GetNextTrack(); } - int localId = TrackIds[(sliceTr->LocalTrackId() >> 24) * mNMaxSingleSliceTracks + (sliceTr->LocalTrackId() & 0xFFFFFF)]; + int localId = mTrackIDs[(sliceTr->LocalTrackId() >> 24) * mNMaxSingleSliceTracks + (sliceTr->LocalTrackId() & 0xFFFFFF)]; if (localId == -1) { continue; } @@ -542,17 +551,15 @@ GPUd() void GPUTPCGMMerger::UnpackSliceGlobal(int nBlocks, int nThreads, int iBl GPUd() void GPUTPCGMMerger::UnpackResetIds(int nBlocks, int nThreads, int iBlock, int iThread, int iSlice) { - int* TrackIds = (int*)mTmpMem; const GPUTPCTracker& trk = GetConstantMem()->tpcTrackers[iSlice]; unsigned int nLocalTracks = Param().rec.mergerReadFromTrackerDirectly ? trk.CommonMemory()->nLocalTracks : mkSlices[iSlice]->NLocalTracks(); for (unsigned int i = iBlock * nThreads + iThread; i < nLocalTracks; i += nBlocks * nThreads) { - TrackIds[iSlice * mNMaxSingleSliceTracks + i] = -1; + mTrackIDs[iSlice * mNMaxSingleSliceTracks + i] = -1; } } GPUd() void GPUTPCGMMerger::RefitSliceTracks(int nBlocks, int nThreads, int iBlock, int iThread, int iSlice) { - int* TrackIds = (int*)mTmpMem; const GPUTPCTracker& trk = GetConstantMem()->tpcTrackers[iSlice]; unsigned int nLocalTracks = Param().rec.mergerReadFromTrackerDirectly ? trk.CommonMemory()->nLocalTracks : mkSlices[iSlice]->NLocalTracks(); @@ -592,7 +599,7 @@ GPUd() void GPUTPCGMMerger::RefitSliceTracks(int nBlocks, int nThreads, int iBlo track.SetGlobalTrackId(0, -1); track.SetGlobalTrackId(1, -1); unsigned int myTrack = CAMath::AtomicAdd(&mMemory->nUnpackedTracks, 1u); - TrackIds[iSlice * mNMaxSingleSliceTracks + sliceTr->LocalTrackId()] = myTrack; + mTrackIDs[iSlice * mNMaxSingleSliceTracks + sliceTr->LocalTrackId()] = myTrack; mSliceTrackInfos[myTrack] = track; } if (!Param().rec.mergerReadFromTrackerDirectly) { @@ -1793,9 +1800,8 @@ GPUd() void GPUTPCGMMerger::SortTracksPrepare(int nBlocks, int nThreads, int iBl GPUd() void GPUTPCGMMerger::PrepareClustersForFit0(int nBlocks, int nThreads, int iBlock, int iThread) { - unsigned int* trackSort = (unsigned int*)mTmpMem; for (unsigned int i = iBlock * nThreads + iThread; i < mMemory->nOutputTracks; i += nBlocks * nThreads) { - trackSort[i] = i; + mTrackSort[i] = i; } } @@ -1841,7 +1847,7 @@ template <> void GPUCA_KRNL_BACKEND_CLASS::runKernelBackendInternal(krnlSetup& _xyz) { GPUDebugTiming timer(mProcessingSettings.debugLevel, nullptr, mInternals->Streams, _xyz, this); - thrust::device_ptr trackSort((unsigned int*)mProcessorsShadow->tpcMerger.TmpMem()); + thrust::device_ptr trackSort((unsigned int*)mProcessorsShadow->tpcMerger.TrackSort()); ThrustVolatileAsyncAllocator alloc(this); thrust::sort(GPUCA_THRUST_NAMESPACE::par(alloc).on(mInternals->Streams[_xyz.x.stream]), trackSort, trackSort + processors()->tpcMerger.NOutputTracks(), GPUTPCGMMergerSortTracksQPt_comp(mProcessorsShadow->tpcMerger.OutputTracks())); } @@ -1873,7 +1879,6 @@ GPUd() void GPUTPCGMMerger::SortTracksQPt(int nBlocks, int nThreads, int iBlock, if (iThread || iBlock) { return; } - unsigned int* trackSort = (unsigned int*)mTmpMem; // Have to duplicate sort comparison: Thrust cannot use the Lambda but OpenCL cannot use the object auto comp = [cmp = mOutputTracks](const int aa, const int bb) { const GPUTPCGMMergedTrack& GPUrestrict() a = cmp[aa]; @@ -1881,20 +1886,18 @@ GPUd() void GPUTPCGMMerger::SortTracksQPt(int nBlocks, int nThreads, int iBlock, return (CAMath::Abs(a.GetParam().GetQPt()) > CAMath::Abs(b.GetParam().GetQPt())); }; - GPUCommonAlgorithm::sortDeviceDynamic(trackSort, trackSort + mMemory->nOutputTracks, comp); + GPUCommonAlgorithm::sortDeviceDynamic(mTrackSort, mTrackSort + mMemory->nOutputTracks, comp); } GPUd() void GPUTPCGMMerger::PrepareClustersForFit1(int nBlocks, int nThreads, int iBlock, int iThread) { - unsigned int* trackSort = (unsigned int*)mTmpMem; - GPUAtomic(unsigned int)* sharedCount = (GPUAtomic(unsigned int)*)(trackSort + CAMath::nextMultipleOf<4>(mMemory->nOutputTracks)); for (unsigned int i = iBlock * nThreads + iThread; i < mMemory->nOutputTracks; i += nBlocks * nThreads) { - mTrackOrderAttach[trackSort[i]] = i; + mTrackOrderAttach[mTrackSort[i]] = i; const GPUTPCGMMergedTrack& trk = mOutputTracks[i]; if (trk.OK()) { for (unsigned int j = 0; j < trk.NClusters(); j++) { mClusterAttachment[mClusters[trk.FirstClusterRef() + j].num] = attachAttached | attachGood; - CAMath::AtomicAdd(&sharedCount[mClusters[trk.FirstClusterRef() + j].num], 1u); + CAMath::AtomicAdd(&mSharedCount[mClusters[trk.FirstClusterRef() + j].num], 1u); } } } @@ -1902,16 +1905,15 @@ GPUd() void GPUTPCGMMerger::PrepareClustersForFit1(int nBlocks, int nThreads, in GPUd() void GPUTPCGMMerger::PrepareClustersForFit2(int nBlocks, int nThreads, int iBlock, int iThread) { - unsigned int* sharedCount = (unsigned int*)mTmpMem + CAMath::nextMultipleOf<4>(mMemory->nOutputTracks); for (unsigned int i = iBlock * nThreads + iThread; i < mMemory->nOutputTrackClusters; i += nBlocks * nThreads) { - if (sharedCount[mClusters[i].num] > 1) { + if (mSharedCount[mClusters[i].num] > 1) { mClusters[i].state |= GPUTPCGMMergedTrackHit::flagShared; } } if (mClusterStateExt) { for (unsigned int i = iBlock * nThreads + iThread; i < mNMaxClusters; i += nBlocks * nThreads) { unsigned char state = GetConstantMem()->ioPtrs.clustersNative->clustersLinear[i].getFlags(); - if (sharedCount[i] > 1) { + if (mSharedCount[i] > 1) { state |= GPUTPCGMMergedTrackHit::flagShared; } mClusterStateExt[i] = state; @@ -1926,9 +1928,8 @@ GPUd() void GPUTPCGMMerger::Finalize0(int nBlocks, int nThreads, int iBlock, int mClusters[i].num = mGlobalClusterIDs[i]; } } else { - int* trkOrderReverse = (int*)mTmpMem; for (unsigned int i = iBlock * nThreads + iThread; i < mMemory->nOutputTracks; i += nThreads * nBlocks) { - trkOrderReverse[mTrackOrderAttach[i]] = i; + mTrackSort[mTrackOrderAttach[i]] = i; } for (unsigned int i = iBlock * nThreads + iThread; i < mMemory->nOutputTrackClusters; i += nThreads * nBlocks) { mClusterAttachment[mClusters[i].num] = 0; // Reset adjacent attachment for attached clusters, set correctly below @@ -1963,10 +1964,9 @@ GPUd() void GPUTPCGMMerger::Finalize1(int nBlocks, int nThreads, int iBlock, int GPUd() void GPUTPCGMMerger::Finalize2(int nBlocks, int nThreads, int iBlock, int iThread) { - int* trkOrderReverse = (int*)mTmpMem; for (unsigned int i = iBlock * nThreads + iThread; i < mNMaxClusters; i += nThreads * nBlocks) { if (mClusterAttachment[i] != 0) { - mClusterAttachment[i] = (mClusterAttachment[i] & attachFlagMask) | trkOrderReverse[mClusterAttachment[i] & attachTrackMask]; + mClusterAttachment[i] = (mClusterAttachment[i] & attachFlagMask) | mTrackSort[mClusterAttachment[i] & attachTrackMask]; } } } diff --git a/GPU/GPUTracking/Merger/GPUTPCGMMerger.h b/GPU/GPUTracking/Merger/GPUTPCGMMerger.h index 2faa7fb37a335..37b82a1f44ce9 100644 --- a/GPU/GPUTracking/Merger/GPUTPCGMMerger.h +++ b/GPU/GPUTracking/Merger/GPUTPCGMMerger.h @@ -81,6 +81,11 @@ class GPUTPCGMMerger : public GPUProcessor unsigned char leg; }; + struct tmpSort { + unsigned int x; + float y; + }; + void InitializeProcessor(); void RegisterMemoryAllocation(); void SetMaxData(const GPUTrackingInOutPointers& io); @@ -120,7 +125,10 @@ class GPUTPCGMMerger : public GPUProcessor GPUhdi() GPUTPCGMLoopData* LoopData() const { return mLoopData; } GPUhdi() memory* Memory() const { return mMemory; } GPUhdi() GPUAtomic(unsigned int) * TmpCounter() { return mMemory->tmpCounter; } - GPUhdi() uint4* TmpMem() { return mTmpMem; } + GPUhdi() uint2* ClusRefTmp() { return mClusRefTmp; } + GPUhdi() unsigned int* TrackSort() { return mTrackSort; } + GPUhdi() tmpSort* TrackSortO2() { return mTrackSortO2; } + GPUhdi() GPUAtomic(unsigned int) * SharedCount() { return mSharedCount; } GPUhdi() gputpcgmmergertypes::GPUTPCGMBorderRange* BorderRange(int i) { return mBorderRange[i]; } GPUhdi() o2::tpc::TrackTPC* OutputTracksTPCO2() { return mOutputTracksTPCO2; } GPUhdi() unsigned int* OutputClusRefsTPCO2() { return mOutputClusRefsTPCO2; } @@ -246,7 +254,11 @@ class GPUTPCGMMerger : public GPUProcessor unsigned int* mTrackOrderAttach; unsigned int* mTrackOrderProcess; unsigned char* mClusterStateExt; - uint4* mTmpMem; + uint2* mClusRefTmp; + int* mTrackIDs; + unsigned int* mTrackSort; + tmpSort* mTrackSortO2; + GPUAtomic(unsigned int) * mSharedCount; GPUTPCGMBorderTrack* mBorderMemory; // memory for border tracks GPUTPCGMBorderTrack* mBorder[2 * NSLICES]; gputpcgmmergertypes::GPUTPCGMBorderRange* mBorderRangeMemory; // memory for border tracks diff --git a/GPU/GPUTracking/Merger/GPUTPCGMO2Output.cxx b/GPU/GPUTracking/Merger/GPUTPCGMO2Output.cxx index 1f771bd8bd00e..4f2fc4ecfadde 100644 --- a/GPU/GPUTracking/Merger/GPUTPCGMO2Output.cxx +++ b/GPU/GPUTracking/Merger/GPUTPCGMO2Output.cxx @@ -31,11 +31,6 @@ using namespace o2::tpc::constants; GPUdi() static constexpr unsigned char getFlagsReject() { return GPUTPCGMMergedTrackHit::flagReject | GPUTPCGMMergedTrackHit::flagNotFit; } GPUdi() static unsigned int getFlagsRequired(const GPUSettingsRec& rec) { return rec.dropSecondaryLegsInOutput ? gputpcgmmergertypes::attachGoodLeg : gputpcgmmergertypes::attachZero; } -struct tmpSort { - unsigned int x; - float y; -}; - template <> GPUdii() void GPUTPCGMO2Output::Thread(int nBlocks, int nThreads, int iBlock, int iThread, GPUsharedref() GPUSharedMemory& smem, processorType& GPUrestrict() merger) { @@ -45,8 +40,8 @@ GPUdii() void GPUTPCGMO2Output::Thread(int nBlocks, i constexpr unsigned char flagsReject = getFlagsReject(); const unsigned int flagsRequired = getFlagsRequired(merger.Param().rec); - tmpSort* trackSort = (tmpSort*)merger.TmpMem(); - uint2* tmpData = (uint2*)(trackSort + merger.NOutputTracks()); + GPUTPCGMMerger::tmpSort* GPUrestrict() trackSort = merger.TrackSortO2(); + uint2* GPUrestrict() tmpData = merger.ClusRefTmp(); for (unsigned int i = get_global_id(0); i < nTracks; i += get_global_size(0)) { unsigned int nCl = 0; for (unsigned int j = 0; j < tracks[i].NClusters(); j++) { @@ -72,14 +67,14 @@ GPUdii() void GPUTPCGMO2Output::Thread(int nBlocks, int if (iThread || iBlock) { return; } - tmpSort* trackSort = (tmpSort*)merger.TmpMem(); + GPUTPCGMMerger::tmpSort* GPUrestrict() trackSort = merger.TrackSortO2(); auto comp = [](const auto& a, const auto& b) { return (a.y > b.y); }; GPUCommonAlgorithm::sortDeviceDynamic(trackSort, trackSort + merger.Memory()->nO2Tracks, comp); } #if (defined(__CUDACC__) || defined(__HIPCC__)) && !defined(GPUCA_GPUCODE_GENRTC) // Specialize GPUTPCGMMergerSortTracks and GPUTPCGMMergerSortTracksQPt struct GPUTPCGMO2OutputSort_comp { - GPUd() bool operator()(const tmpSort& a, const tmpSort& b) + GPUd() bool operator()(const GPUTPCGMMerger::tmpSort& a, const GPUTPCGMMerger::tmpSort& b) { return (a.y > b.y); } @@ -89,7 +84,7 @@ template <> void GPUCA_KRNL_BACKEND_CLASS::runKernelBackendInternal(krnlSetup& _xyz) { GPUDebugTiming timer(mProcessingSettings.debugLevel, nullptr, mInternals->Streams, _xyz, this); - thrust::device_ptr trackSort((tmpSort*)mProcessorsShadow->tpcMerger.TmpMem()); + thrust::device_ptr trackSort(mProcessorsShadow->tpcMerger.TrackSortO2()); ThrustVolatileAsyncAllocator alloc(this); thrust::sort(GPUCA_THRUST_NAMESPACE::par(alloc).on(mInternals->Streams[_xyz.x.stream]), trackSort, trackSort + processors()->tpcMerger.NOutputTracksTPCO2(), GPUTPCGMO2OutputSort_comp()); } @@ -107,8 +102,8 @@ GPUdii() void GPUTPCGMO2Output::Thread(int nBlocks, in TrackTPC* outputTracks = merger.OutputTracksTPCO2(); unsigned int* clusRefs = merger.OutputClusRefsTPCO2(); - tmpSort* trackSort = (tmpSort*)merger.TmpMem(); - uint2* tmpData = (uint2*)(trackSort + merger.NOutputTracks()); + GPUTPCGMMerger::tmpSort* GPUrestrict() trackSort = merger.TrackSortO2(); + uint2* GPUrestrict() tmpData = merger.ClusRefTmp(); for (int iTmp = get_global_id(0); iTmp < nTracks; iTmp += get_global_size(0)) { auto& oTrack = outputTracks[iTmp]; From 32d01b9908cc8fc0985e5257fabcf44af1c41006 Mon Sep 17 00:00:00 2001 From: David Rohr Date: Tue, 2 Mar 2021 20:47:21 +0100 Subject: [PATCH 09/20] GPU: More memory reuise in TPC track merging --- GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx | 10 +++++----- 1 file changed, 5 insertions(+), 5 deletions(-) diff --git a/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx b/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx index da6cee561adb0..4d1de3c683c12 100644 --- a/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx +++ b/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx @@ -249,7 +249,8 @@ void* GPUTPCGMMerger::SetPointersMerger(void* mem) computePointerWithAlignment(mem, mGlobalClusterIDs, mNMaxOutputTrackClusters); } - computePointerWithAlignment(mem, mBorderMemory, 2 * mNMaxSliceTracks); + void* memBase = mem; + computePointerWithAlignment(mem, mBorderMemory, 2 * mNMaxSliceTracks); // MergeBorders & Resolve computePointerWithAlignment(mem, mBorderRangeMemory, 2 * mNMaxSliceTracks); int nTracks = 0; for (int iSlice = 0; iSlice < NSLICES; iSlice++) { @@ -259,14 +260,13 @@ void* GPUTPCGMMerger::SetPointersMerger(void* mem) mBorderRange[iSlice] = mBorderRangeMemory + 2 * nTracks; nTracks += n; } - computePointerWithAlignment(mem, mTrackLinks, mNMaxSliceTracks); computePointerWithAlignment(mem, mTrackCCRoots, mNMaxSliceTracks); - - void* memBase = mem; + void* memMax = mem; + mem = memBase; computePointerWithAlignment(mem, mTrackSortO2, mNMaxTracks); // GPUTPCGMO2Output::prepare - GPUTPCGMO2Output::sort - GPUTPCGMO2Output::output computePointerWithAlignment(mem, mClusRefTmp, mNMaxTracks); - void* memMax = mem; + memMax = std::max(mem, memMax); mem = memBase; computePointerWithAlignment(mem, mTrackIDs, mNMaxTracks); // UnpackResetIds - RefitSliceTracks - UnpackSliceGlobal memMax = std::max(mem, memMax); From 9f4aed6cbfd98ea12033cee6ba57d3f9a9c595f4 Mon Sep 17 00:00:00 2001 From: David Rohr Date: Tue, 2 Mar 2021 22:02:02 +0100 Subject: [PATCH 10/20] GPU: Workaround for GCC compiler bug --- GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx b/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx index 4d1de3c683c12..ee22f582d39e3 100644 --- a/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx +++ b/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx @@ -266,18 +266,18 @@ void* GPUTPCGMMerger::SetPointersMerger(void* mem) mem = memBase; computePointerWithAlignment(mem, mTrackSortO2, mNMaxTracks); // GPUTPCGMO2Output::prepare - GPUTPCGMO2Output::sort - GPUTPCGMO2Output::output computePointerWithAlignment(mem, mClusRefTmp, mNMaxTracks); - memMax = std::max(mem, memMax); + memMax = (void*)std::max((size_t)mem, (size_t)memMax); // BUG: Casts needed to workaround GCC miscompilation mem = memBase; computePointerWithAlignment(mem, mTrackIDs, mNMaxTracks); // UnpackResetIds - RefitSliceTracks - UnpackSliceGlobal - memMax = std::max(mem, memMax); + memMax = (void*)std::max((size_t)mem, (size_t)memMax); mem = memBase; computePointerWithAlignment(mem, mTrackSort, mNMaxTracks); // PrepareClustersForFit0 - SortTracksQPt - PrepareClustersForFit1 - PrepareClustersForFit1 / Finalize0 - Finalize2 computePointerWithAlignment(mem, mSharedCount, mNMaxClusters); - memMax = std::max(mem, memMax); + memMax = (void*)std::max((size_t)mem, (size_t)memMax); mem = memBase; computePointerWithAlignment(mem, mLoopData, mNMaxTracks); // GPUTPCGMMergerTrackFit - GPUTPCGMMergerFollowLoopers computePointerWithAlignment(mem, mRetryRefitIds, mNMaxTracks); - memMax = std::max(mem, memMax); + memMax = (void*)std::max((size_t)mem, (size_t)memMax); return memMax; } From 6d652d6e808fce60f56d4ae8d3a15d9296f48cc6 Mon Sep 17 00:00:00 2001 From: David Rohr Date: Tue, 2 Mar 2021 23:17:31 +0100 Subject: [PATCH 11/20] GPU: Add some notes --- GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx | 4 ++-- GPU/GPUTracking/Merger/GPUTPCGMMerger.h | 2 +- 2 files changed, 3 insertions(+), 3 deletions(-) diff --git a/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx b/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx index ee22f582d39e3..95ba44e3b0427 100644 --- a/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx +++ b/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx @@ -275,8 +275,8 @@ void* GPUTPCGMMerger::SetPointersMerger(void* mem) computePointerWithAlignment(mem, mSharedCount, mNMaxClusters); memMax = (void*)std::max((size_t)mem, (size_t)memMax); mem = memBase; - computePointerWithAlignment(mem, mLoopData, mNMaxTracks); // GPUTPCGMMergerTrackFit - GPUTPCGMMergerFollowLoopers - computePointerWithAlignment(mem, mRetryRefitIds, mNMaxTracks); + computePointerWithAlignment(mem, mLoopData, mNMaxTracks); // GPUTPCGMMergerTrackFit - GPUTPCGMMergerFollowLoopers + computePointerWithAlignment(mem, mRetryRefitIds, mNMaxTracks); // Reducing mNMaxTracks for mLoopData / mRetryRefitIds does not save memory, since the other parts are larger anyway memMax = (void*)std::max((size_t)mem, (size_t)memMax); return memMax; } diff --git a/GPU/GPUTracking/Merger/GPUTPCGMMerger.h b/GPU/GPUTracking/Merger/GPUTPCGMMerger.h index 37b82a1f44ce9..8cc2ba55a44fa 100644 --- a/GPU/GPUTracking/Merger/GPUTPCGMMerger.h +++ b/GPU/GPUTracking/Merger/GPUTPCGMMerger.h @@ -258,7 +258,7 @@ class GPUTPCGMMerger : public GPUProcessor int* mTrackIDs; unsigned int* mTrackSort; tmpSort* mTrackSortO2; - GPUAtomic(unsigned int) * mSharedCount; + GPUAtomic(unsigned int) * mSharedCount; // Must be unsigned int unfortunately for atomic support GPUTPCGMBorderTrack* mBorderMemory; // memory for border tracks GPUTPCGMBorderTrack* mBorder[2 * NSLICES]; gputpcgmmergertypes::GPUTPCGMBorderRange* mBorderRangeMemory; // memory for border tracks From 96272a03080a92dfc602a188334cf290b6b58552 Mon Sep 17 00:00:00 2001 From: David Rohr Date: Wed, 3 Mar 2021 00:00:09 +0100 Subject: [PATCH 12/20] GPU: Enforce correct alignment for preallocated buffers in standalone benchmark --- GPU/GPUTracking/Benchmark/standalone.cxx | 12 ++++++++---- 1 file changed, 8 insertions(+), 4 deletions(-) diff --git a/GPU/GPUTracking/Benchmark/standalone.cxx b/GPU/GPUTracking/Benchmark/standalone.cxx index 0fe173e22c5a9..2c0ea3af870d3 100644 --- a/GPU/GPUTracking/Benchmark/standalone.cxx +++ b/GPU/GPUTracking/Benchmark/standalone.cxx @@ -82,7 +82,11 @@ GPUChainTracking *chainTracking, *chainTrackingAsync, *chainTrackingPipeline; #ifdef HAVE_O2HEADERS GPUChainITS *chainITS, *chainITSAsync, *chainITSPipeline; #endif -std::unique_ptr outputmemory, outputmemoryPipeline, inputmemory; +void unique_ptr_aligned_delete(char* v) +{ + operator delete(v GPUCA_OPERATOR_NEW_ALIGNMENT); +} +std::unique_ptr outputmemory(nullptr, unique_ptr_aligned_delete), outputmemoryPipeline(nullptr, unique_ptr_aligned_delete), inputmemory(nullptr, unique_ptr_aligned_delete); std::unique_ptr eventDisplay; std::unique_ptr tf; int nEventsInDirectory = 0; @@ -237,20 +241,20 @@ int ReadConfiguration(int argc, char** argv) if (configStandalone.outputcontrolmem) { bool forceEmptyMemory = getenv("LD_PRELOAD") && strstr(getenv("LD_PRELOAD"), "valgrind") != nullptr; - outputmemory.reset(new char[configStandalone.outputcontrolmem]); + outputmemory.reset((char*)operator new(configStandalone.outputcontrolmem GPUCA_OPERATOR_NEW_ALIGNMENT)); if (forceEmptyMemory) { printf("Valgrind detected, emptying GPU output memory to avoid false positive undefined reads"); memset(outputmemory.get(), 0, configStandalone.outputcontrolmem); } if (configStandalone.proc.doublePipeline) { - outputmemoryPipeline.reset(new char[configStandalone.outputcontrolmem]); + outputmemoryPipeline.reset((char*)operator new(configStandalone.outputcontrolmem GPUCA_OPERATOR_NEW_ALIGNMENT)); if (forceEmptyMemory) { memset(outputmemoryPipeline.get(), 0, configStandalone.outputcontrolmem); } } } if (configStandalone.inputcontrolmem) { - inputmemory.reset(new char[configStandalone.inputcontrolmem]); + inputmemory.reset((char*)operator new(configStandalone.inputcontrolmem GPUCA_OPERATOR_NEW_ALIGNMENT)); } #if !(defined(CUDA_ENABLED) || defined(OPENCL1_ENABLED) || defined(HIP_ENABLED)) From 8cea57bf23eb89731574a6e2a0e1bb8694859ac1 Mon Sep 17 00:00:00 2001 From: David Rohr Date: Wed, 3 Mar 2021 09:37:45 +0100 Subject: [PATCH 13/20] GPU: Allocate dEdx memory only when needed --- GPU/GPUTracking/Global/GPUChainTrackingMerger.cxx | 3 +++ GPU/GPUTracking/Merger/GPUTPCGMMergedTrack.h | 12 +----------- GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx | 3 +++ GPU/GPUTracking/Merger/GPUTPCGMMerger.h | 13 ++++++------- GPU/GPUTracking/Merger/GPUTPCGMMergerDump.cxx | 3 ++- GPU/GPUTracking/Merger/GPUTPCGMO2Output.cxx | 5 ++++- GPU/GPUTracking/Merger/GPUTPCGMTrackParam.cxx | 14 +++++++------- GPU/GPUTracking/Merger/GPUTPCGMTrackParam.h | 2 +- 8 files changed, 27 insertions(+), 28 deletions(-) diff --git a/GPU/GPUTracking/Global/GPUChainTrackingMerger.cxx b/GPU/GPUTracking/Global/GPUChainTrackingMerger.cxx index 55ec254b4668c..4fab18d788ecd 100644 --- a/GPU/GPUTracking/Global/GPUChainTrackingMerger.cxx +++ b/GPU/GPUTracking/Global/GPUChainTrackingMerger.cxx @@ -241,6 +241,9 @@ int GPUChainTracking::RunTPCTrackingMerger(bool synchronizeOutput) TransferMemoryResourceLinkToHost(RecoStep::TPCMerging, Merger.MemoryResOutput(), mRec->NStreams() - 2, nullptr, &mEvents->single); } else if (GetProcessingSettings().keepDisplayMemory || GetProcessingSettings().createO2Output <= 1) { GPUMemCpy(RecoStep::TPCMerging, Merger.OutputTracks(), MergerShadowAll.OutputTracks(), Merger.NOutputTracks() * sizeof(*Merger.OutputTracks()), mRec->NStreams() - 2, 0, nullptr, &mEvents->single); + if (param().par.dodEdx) { + GPUMemCpy(RecoStep::TPCMerging, Merger.OutputTracksdEdx(), MergerShadowAll.OutputTracksdEdx(), Merger.NOutputTracks() * sizeof(*Merger.OutputTracksdEdx()), mRec->NStreams() - 2, 0, nullptr, &mEvents->single); + } GPUMemCpy(RecoStep::TPCMerging, Merger.Clusters(), MergerShadowAll.Clusters(), Merger.NOutputTrackClusters() * sizeof(*Merger.Clusters()), mRec->NStreams() - 2, 0); if (param().par.earlyTpcTransform) { GPUMemCpy(RecoStep::TPCMerging, Merger.ClustersXYZ(), MergerShadowAll.ClustersXYZ(), Merger.NOutputTrackClusters() * sizeof(*Merger.ClustersXYZ()), mRec->NStreams() - 2, 0); diff --git a/GPU/GPUTracking/Merger/GPUTPCGMMergedTrack.h b/GPU/GPUTracking/Merger/GPUTPCGMMergedTrack.h index 1a4479fc74d38..eeed5d24cc82f 100644 --- a/GPU/GPUTracking/Merger/GPUTPCGMMergedTrack.h +++ b/GPU/GPUTracking/Merger/GPUTPCGMMergedTrack.h @@ -16,7 +16,6 @@ #include "GPUTPCGMTrackParam.h" #include "GPUTPCGMMergedTrackHit.h" -#include "GPUdEdxInfo.h" namespace GPUCA_NAMESPACE { @@ -105,20 +104,11 @@ class GPUTPCGMMergedTrack GPUd() unsigned char Legs() const { return mLegs; } GPUd() const GPUTPCGMTrackParam::GPUTPCOuterParam& OuterParam() const { return mOuterParam; } - GPUd() GPUTPCGMTrackParam::GPUTPCOuterParam& OuterParam() - { - return mOuterParam; - } - GPUd() const GPUdEdxInfo& dEdxInfo() const { return mdEdxInfo; } - GPUd() GPUdEdxInfo& dEdxInfo() - { - return mdEdxInfo; - } + GPUd() GPUTPCGMTrackParam::GPUTPCOuterParam& OuterParam() { return mOuterParam; } private: GPUTPCGMTrackParam mParam; //* fitted track parameters GPUTPCGMTrackParam::GPUTPCOuterParam mOuterParam; //* outer param - GPUdEdxInfo mdEdxInfo; //* dEdx information float mAlpha; //* alpha angle float mLastX; //* outer X diff --git a/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx b/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx index 95ba44e3b0427..a197526f2e95f 100644 --- a/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx +++ b/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx @@ -307,6 +307,9 @@ void* GPUTPCGMMerger::SetPointersRefitScratch2(void* mem) void* GPUTPCGMMerger::SetPointersOutput(void* mem) { computePointerWithAlignment(mem, mOutputTracks, mNMaxTracks); + if (mRec->GetParam().par.dodEdx) { + computePointerWithAlignment(mem, mOutputTracksdEdx, mNMaxTracks); + } computePointerWithAlignment(mem, mClusters, mNMaxOutputTrackClusters); if (mRec->GetParam().par.earlyTpcTransform) { computePointerWithAlignment(mem, mClustersXYZ, mNMaxOutputTrackClusters); diff --git a/GPU/GPUTracking/Merger/GPUTPCGMMerger.h b/GPU/GPUTracking/Merger/GPUTPCGMMerger.h index 8cc2ba55a44fa..44c3599387176 100644 --- a/GPU/GPUTracking/Merger/GPUTPCGMMerger.h +++ b/GPU/GPUTracking/Merger/GPUTPCGMMerger.h @@ -101,13 +101,11 @@ class GPUTPCGMMerger : public GPUProcessor void SetSliceData(int index, const GPUTPCSliceOutput* sliceData) { mkSlices[index] = sliceData; } - GPUhd() int NOutputTracks() const { return mMemory->nOutputTracks; } - GPUhd() const GPUTPCGMMergedTrack* OutputTracks() const { return mOutputTracks; } - GPUhd() GPUTPCGMMergedTrack* OutputTracks() - { - return mOutputTracks; - } - + GPUhdi() int NOutputTracks() const { return mMemory->nOutputTracks; } + GPUhdi() const GPUTPCGMMergedTrack* OutputTracks() const { return mOutputTracks; } + GPUhdi() GPUTPCGMMergedTrack* OutputTracks() { return mOutputTracks; } + GPUhdi() const GPUdEdxInfo* OutputTracksdEdx() const { return mOutputTracksdEdx; } + GPUhdi() GPUdEdxInfo* OutputTracksdEdx() { return mOutputTracksdEdx; } GPUhdi() unsigned int NClusters() const { return mNClusters; } GPUhdi() unsigned int NMaxClusters() const { return mNMaxClusters; } GPUhdi() unsigned int NMaxTracks() const { return mNMaxTracks; } @@ -241,6 +239,7 @@ class GPUTPCGMMerger : public GPUProcessor int mNClusters; // Total number of incoming clusters (from slice tracks) GPUTPCGMMergedTrack* mOutputTracks; //* array of output merged tracks + GPUdEdxInfo* mOutputTracksdEdx; //* dEdx information GPUTPCGMSliceTrack* mSliceTrackInfos; //* additional information for slice tracks int* mSliceTrackInfoIndex; GPUTPCGMMergedTrackHit* mClusters; diff --git a/GPU/GPUTracking/Merger/GPUTPCGMMergerDump.cxx b/GPU/GPUTracking/Merger/GPUTPCGMMergerDump.cxx index d9c9e0c5cec23..3474e33c32b84 100644 --- a/GPU/GPUTracking/Merger/GPUTPCGMMergerDump.cxx +++ b/GPU/GPUTracking/Merger/GPUTPCGMMergerDump.cxx @@ -154,11 +154,12 @@ void GPUTPCGMMerger::DumpRefit(std::ostream& out) out << "\nTPC Merger Refit\n"; for (unsigned int i = 0; i < mMemory->nOutputTracks; i++) { const auto& trk = mOutputTracks[trackOrder[i]]; + const auto& trkdEdx = mOutputTracksdEdx[trackOrder[i]]; const auto& p = trk.GetParam(); const auto& po = trk.OuterParam(); out << " Track " << i << ": OK " << trk.OK() << " Alpha " << trk.GetAlpha() << " X " << p.GetX() << " Y " << p.GetY() << " Z " << p.GetZ() << " SPhi " << p.GetSinPhi() << " Tgl " << p.GetDzDs() << " QPt " << p.GetQPt() << " NCl " << trk.NClusters() << " / " << trk.NClustersFitted() << " Cov " << p.GetErr2Y() << "/" << p.GetErr2Z() #ifdef HAVE_O2HEADERS - << " dEdx " << trk.dEdxInfo().dEdxTotTPC << "/" << trk.dEdxInfo().dEdxMaxTPC + << " dEdx " << trkdEdx.dEdxTotTPC << "/" << trkdEdx.dEdxMaxTPC #endif << " Outer " << po.P[0] << "/" << po.P[1] << "/" << po.P[2] << "/" << po.P[3] << "/" << po.P[4] << "\n"; } diff --git a/GPU/GPUTracking/Merger/GPUTPCGMO2Output.cxx b/GPU/GPUTracking/Merger/GPUTPCGMO2Output.cxx index 4f2fc4ecfadde..2354076773341 100644 --- a/GPU/GPUTracking/Merger/GPUTPCGMO2Output.cxx +++ b/GPU/GPUTracking/Merger/GPUTPCGMO2Output.cxx @@ -95,6 +95,7 @@ GPUdii() void GPUTPCGMO2Output::Thread(int nBlocks, in { constexpr float MinDelta = 0.1; const GPUTPCGMMergedTrack* tracks = merger.OutputTracks(); + GPUdEdxInfo* tracksdEdx = merger.OutputTracksdEdx(); const int nTracks = merger.NOutputTracksTPCO2(); const GPUTPCGMMergedTrackHit* trackClusters = merger.Clusters(); constexpr unsigned char flagsReject = getFlagsReject(); @@ -119,7 +120,9 @@ GPUdii() void GPUTPCGMO2Output::Thread(int nBlocks, in oTrack.setChi2(tracks[i].GetParam().GetChi2()); auto& outerPar = tracks[i].OuterParam(); - oTrack.setdEdx(tracks[i].dEdxInfo()); + if (merger.Param().par.dodEdx) { + oTrack.setdEdx(tracksdEdx[i]); + } oTrack.setOuterParam(o2::track::TrackParCov( outerPar.X, outerPar.alpha, {outerPar.P[0], outerPar.P[1], outerPar.P[2], outerPar.P[3], outerPar.P[4]}, diff --git a/GPU/GPUTracking/Merger/GPUTPCGMTrackParam.cxx b/GPU/GPUTracking/Merger/GPUTPCGMTrackParam.cxx index f642cae803cb0..cc980905e7612 100644 --- a/GPU/GPUTracking/Merger/GPUTPCGMTrackParam.cxx +++ b/GPU/GPUTracking/Merger/GPUTPCGMTrackParam.cxx @@ -54,7 +54,7 @@ using namespace o2::tpc; static constexpr float kDeg2Rad = M_PI / 180.f; static constexpr float kSectAngle = 2 * M_PI / 18.f; -GPUd() bool GPUTPCGMTrackParam::Fit(const GPUTPCGMMerger* GPUrestrict() merger, int iTrk, GPUTPCGMMergedTrackHit* GPUrestrict() clusters, GPUTPCGMMergedTrackHitXYZ* GPUrestrict() clustersXYZ, int& GPUrestrict() N, int& GPUrestrict() NTolerated, float& GPUrestrict() Alpha, int attempt, float maxSinPhi, GPUTPCOuterParam* GPUrestrict() outerParam, GPUdEdxInfo* GPUrestrict() dEdxOut) +GPUd() bool GPUTPCGMTrackParam::Fit(GPUTPCGMMerger* GPUrestrict() merger, int iTrk, GPUTPCGMMergedTrackHit* GPUrestrict() clusters, GPUTPCGMMergedTrackHitXYZ* GPUrestrict() clustersXYZ, int& GPUrestrict() N, int& GPUrestrict() NTolerated, float& GPUrestrict() Alpha, int attempt, float maxSinPhi, GPUTPCOuterParam* GPUrestrict() outerParam) { const GPUParam& GPUrestrict() param = merger->Param(); @@ -196,7 +196,7 @@ GPUd() bool GPUTPCGMTrackParam::Fit(const GPUTPCGMMerger* GPUrestrict() merger, continue; } } else if (allowModification && lastRow != 255 && CAMath::Abs(clusters[ihit].row - lastRow) > 1) { - if (dEdxOut && iWay == nWays - 1 && clusters[ihit].row == lastRow - 2 && clusters[ihit].leg == clusters[maxN - 1].leg) { + if (merger->Param().par.dodEdx && iWay == nWays - 1 && clusters[ihit].row == lastRow - 2 && clusters[ihit].leg == clusters[maxN - 1].leg) { dEdx.fillSubThreshold(lastRow - 1, param); } AttachClustersPropagate(merger, clusters[ihit].slice, lastRow, clusters[ihit].row, iTrk, clusters[ihit].leg == clusters[maxN - 1].leg, prop, inFlyDirection); @@ -296,7 +296,7 @@ GPUd() bool GPUTPCGMTrackParam::Fit(const GPUTPCGMMerger* GPUrestrict() merger, CADEBUG(printf("Reinit linearization\n")); prop.SetTrack(this, prop.GetAlpha()); } - if (dEdxOut && iWay == nWays - 1 && clusters[ihit].leg == clusters[maxN - 1].leg) { + if (merger->Param().par.dodEdx && iWay == nWays - 1 && clusters[ihit].leg == clusters[maxN - 1].leg) { float qtot, qmax; if (merger->GetConstantMem()->ioPtrs.clustersNative == nullptr) { qtot = clustersXYZ[ihit].amp; @@ -331,8 +331,8 @@ GPUd() bool GPUTPCGMTrackParam::Fit(const GPUTPCGMMerger* GPUrestrict() merger, } // TODO: we have looping tracks here with 0 accepted clusters in the primary leg. In that case we should refit the track using only the primary leg. - if (dEdxOut) { - dEdx.computedEdx(*dEdxOut, param); + if (merger->Param().par.dodEdx) { + dEdx.computedEdx(merger->OutputTracksdEdx()[iTrk], param); } Alpha = prop.GetAlpha(); MoveToReference(prop, param, Alpha); @@ -982,7 +982,7 @@ GPUd() void GPUTPCGMTrackParam::RefitTrack(GPUTPCGMMergedTrack& GPUrestrict() tr GPUTPCGMTrackParam t = track.Param(); float Alpha = track.Alpha(); CADEBUG(int nTrackHitsOld = nTrackHits; float ptOld = t.QPt()); - bool ok = t.Fit(merger, iTrk, merger->Clusters() + track.FirstClusterRef(), merger->ClustersXYZ() + track.FirstClusterRef(), nTrackHits, NTolerated, Alpha, attempt, GPUCA_MAX_SIN_PHI, &track.OuterParam(), merger->Param().par.dodEdx ? &track.dEdxInfo() : nullptr); + bool ok = t.Fit(merger, iTrk, merger->Clusters() + track.FirstClusterRef(), merger->ClustersXYZ() + track.FirstClusterRef(), nTrackHits, NTolerated, Alpha, attempt, GPUCA_MAX_SIN_PHI, &track.OuterParam()); CADEBUG(printf("Finished Fit Track %d\n", iTrk)); CADEBUG(printf("OUTPUT hits %d -> %d+%d = %d, QPt %f -> %f, SP %f, ok %d chi2 %f chi2ndf %f\n", nTrackHitsOld, nTrackHits, NTolerated, nTrackHits + NTolerated, ptOld, t.QPt(), t.SinPhi(), (int)ok, t.Chi2(), t.Chi2() / CAMath::Max(1, nTrackHits))); @@ -996,7 +996,7 @@ GPUd() void GPUTPCGMTrackParam::RefitTrack(GPUTPCGMMergedTrack& GPUrestrict() tr NTolerated = 0; // Clusters not fit but tollerated for track length cut t = track.Param(); Alpha = track.Alpha(); - ok = t.Fit(merger, iTrk, merger->Clusters() + track.FirstClusterRef(), merger->ClustersXYZ() + track.FirstClusterRef(), nTrackHits, NTolerated, Alpha, 1, GPUCA_MAX_SIN_PHI, &track.OuterParam(), merger->Param().par.dodEdx ? &track.dEdxInfo() : nullptr); + ok = t.Fit(merger, iTrk, merger->Clusters() + track.FirstClusterRef(), merger->ClustersXYZ() + track.FirstClusterRef(), nTrackHits, NTolerated, Alpha, 1, GPUCA_MAX_SIN_PHI, &track.OuterParam()); } else { unsigned int nRefit = CAMath::AtomicAdd(&merger->Memory()->nRetryRefit, 1u); merger->RetryRefitIds()[nRefit] = iTrk; diff --git a/GPU/GPUTracking/Merger/GPUTPCGMTrackParam.h b/GPU/GPUTracking/Merger/GPUTPCGMTrackParam.h index 24be2916ed69f..96f95a01876b3 100644 --- a/GPU/GPUTracking/Merger/GPUTPCGMTrackParam.h +++ b/GPU/GPUTracking/Merger/GPUTPCGMTrackParam.h @@ -147,7 +147,7 @@ class GPUTPCGMTrackParam GPUd() bool CheckNumericalQuality(float overrideCovYY = -1.f) const; GPUd() bool CheckCov() const; - GPUd() bool Fit(const GPUTPCGMMerger* merger, int iTrk, GPUTPCGMMergedTrackHit* clusters, GPUTPCGMMergedTrackHitXYZ* clustersXYZ, int& N, int& NTolerated, float& Alpha, int attempt = 0, float maxSinPhi = GPUCA_MAX_SIN_PHI, GPUTPCOuterParam* outerParam = nullptr, GPUdEdxInfo* dEdxOut = nullptr); + GPUd() bool Fit(GPUTPCGMMerger* merger, int iTrk, GPUTPCGMMergedTrackHit* clusters, GPUTPCGMMergedTrackHitXYZ* clustersXYZ, int& N, int& NTolerated, float& Alpha, int attempt = 0, float maxSinPhi = GPUCA_MAX_SIN_PHI, GPUTPCOuterParam* outerParam = nullptr); GPUd() void MoveToReference(GPUTPCGMPropagator& prop, const GPUParam& param, float& alpha); GPUd() void MirrorTo(GPUTPCGMPropagator& prop, float toY, float toZ, bool inFlyDirection, const GPUParam& param, unsigned char row, unsigned char clusterState, bool mirrorParameters); GPUd() int MergeDoubleRowClusters(int& ihit, int wayDirection, GPUTPCGMMergedTrackHit* clusters, GPUTPCGMMergedTrackHitXYZ* clustersXYZ, const GPUTPCGMMerger* merger, GPUTPCGMPropagator& prop, float& xx, float& yy, float& zz, int maxN, float clAlpha, unsigned char& clusterState, bool rejectChi2); From f437a4e18af5b88e9b0537a8f6594bae55436afb Mon Sep 17 00:00:00 2001 From: David Rohr Date: Wed, 3 Mar 2021 10:40:43 +0100 Subject: [PATCH 14/20] GPU: Add utility class to create int tags from const strings --- GPU/GPUTracking/utils/strtag.h | 38 ++++++++++++++++++++++++++++++++++ 1 file changed, 38 insertions(+) create mode 100644 GPU/GPUTracking/utils/strtag.h diff --git a/GPU/GPUTracking/utils/strtag.h b/GPU/GPUTracking/utils/strtag.h new file mode 100644 index 0000000000000..bf8a92ee2e0ff --- /dev/null +++ b/GPU/GPUTracking/utils/strtag.h @@ -0,0 +1,38 @@ +// Copyright CERN and copyright holders of ALICE O2. This software is +// distributed under the terms of the GNU General Public License v3 (GPL +// Version 3), copied verbatim in the file "COPYING". +// +// See http://alice-o2.web.cern.ch/license for full licensing information. +// +// 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. + +/// \file strtag.h +/// \author David Rohr + +#ifndef STRTAG_H +#define STRTAG_H + +#include +#include + +template +constexpr T qStr2Tag(const char* str) +{ + if (strlen(str) != sizeof(T)) { + throw std::runtime_error("Invalid tag length"); + } + return *(T*)str; +} + +template +std::string qTag2Str(const T tag) +{ + T str[2]; + str[0] = tag; + str[1] = 0; + return std::string((const char*)str); +} + +#endif From 940c34626d39a04bffdbe3f76058d115eb5c3e2b Mon Sep 17 00:00:00 2001 From: David Rohr Date: Wed, 3 Mar 2021 10:53:04 +0100 Subject: [PATCH 15/20] GPU: Protect non-persistent memory stack with tags to avoid accidental inconsistent usage --- GPU/GPUTracking/Base/GPUReconstruction.cxx | 14 +++++++++----- GPU/GPUTracking/Base/GPUReconstruction.h | 6 +++--- GPU/GPUTracking/Global/GPUChainTracking.cxx | 5 +++-- .../Global/GPUChainTrackingClusterizer.cxx | 6 ++++-- .../Global/GPUChainTrackingCompression.cxx | 5 +++-- GPU/GPUTracking/Global/GPUChainTrackingMerger.cxx | 5 +++-- .../Global/GPUChainTrackingSliceTracker.cxx | 5 +++-- GPU/GPUTracking/Global/GPUChainTrackingTRD.cxx | 5 +++-- .../Global/GPUChainTrackingTransformation.cxx | 5 +++-- 9 files changed, 34 insertions(+), 22 deletions(-) diff --git a/GPU/GPUTracking/Base/GPUReconstruction.cxx b/GPU/GPUTracking/Base/GPUReconstruction.cxx index 1db577b8b31f2..74f42ce24568d 100644 --- a/GPU/GPUTracking/Base/GPUReconstruction.cxx +++ b/GPU/GPUTracking/Base/GPUReconstruction.cxx @@ -42,6 +42,7 @@ #include "GPUChain.h" #include "GPUMemorySizeScalers.h" +#include "utils/strtag.h" #define GPUCA_LOGGING_PRINTF #include "GPULogging.h" @@ -722,18 +723,21 @@ void GPUReconstruction::ReturnVolatileDeviceMemory() } } -void GPUReconstruction::PushNonPersistentMemory() +void GPUReconstruction::PushNonPersistentMemory(unsigned long tag) { - mNonPersistentMemoryStack.emplace_back(mHostMemoryPoolEnd, mDeviceMemoryPoolEnd, mNonPersistentIndividualAllocations.size()); + mNonPersistentMemoryStack.emplace_back(mHostMemoryPoolEnd, mDeviceMemoryPoolEnd, mNonPersistentIndividualAllocations.size(), tag); } -void GPUReconstruction::PopNonPersistentMemory(RecoStep step) +void GPUReconstruction::PopNonPersistentMemory(RecoStep step, unsigned long tag) { + if (tag != 0 && std::get<3>(mNonPersistentMemoryStack.back()) != tag) { + GPUFatal("Tag mismatch when poping non persistent memory from stack : pop %s vs on stack %s", qTag2Str(tag).c_str(), qTag2Str(std::get<3>(mNonPersistentMemoryStack.back())).c_str()); + } if ((mProcessingSettings.debugLevel >= 3 || mProcessingSettings.allocDebugLevel) && (IsGPU() || mProcessingSettings.forceHostMemoryPoolSize)) { if (IsGPU()) { - printf("Allocated Device memory after %30s: %'13lld (non temporary %'13lld, blocked %'13lld)\n", GPUDataTypes::RECO_STEP_NAMES[getRecoStepNum(step, true)], ptrDiff(mDeviceMemoryPool, mDeviceMemoryBase) + ptrDiff((char*)mDeviceMemoryBase + mDeviceMemorySize, mDeviceMemoryPoolEnd), ptrDiff(mDeviceMemoryPool, mDeviceMemoryBase), mDeviceMemoryPoolBlocked == nullptr ? 0ll : ptrDiff((char*)mDeviceMemoryBase + mDeviceMemorySize, mDeviceMemoryPoolBlocked)); + printf("Allocated Device memory after %30s (%8s): %'13lld (non temporary %'13lld, blocked %'13lld)\n", GPUDataTypes::RECO_STEP_NAMES[getRecoStepNum(step, true)], qTag2Str(std::get<3>(mNonPersistentMemoryStack.back())).c_str(), ptrDiff(mDeviceMemoryPool, mDeviceMemoryBase) + ptrDiff((char*)mDeviceMemoryBase + mDeviceMemorySize, mDeviceMemoryPoolEnd), ptrDiff(mDeviceMemoryPool, mDeviceMemoryBase), mDeviceMemoryPoolBlocked == nullptr ? 0ll : ptrDiff((char*)mDeviceMemoryBase + mDeviceMemorySize, mDeviceMemoryPoolBlocked)); } - printf("Allocated Host memory after %30s: %'13lld (non temporary %'13lld, blocked %'13lld)\n", GPUDataTypes::RECO_STEP_NAMES[getRecoStepNum(step, true)], ptrDiff(mHostMemoryPool, mHostMemoryBase) + ptrDiff((char*)mHostMemoryBase + mHostMemorySize, mHostMemoryPoolEnd), ptrDiff(mHostMemoryPool, mHostMemoryBase), mHostMemoryPoolBlocked == nullptr ? 0ll : ptrDiff((char*)mHostMemoryBase + mHostMemorySize, mHostMemoryPoolBlocked)); + printf("Allocated Host memory after %30s (%8s): %'13lld (non temporary %'13lld, blocked %'13lld)\n", GPUDataTypes::RECO_STEP_NAMES[getRecoStepNum(step, true)], qTag2Str(std::get<3>(mNonPersistentMemoryStack.back())).c_str(), ptrDiff(mHostMemoryPool, mHostMemoryBase) + ptrDiff((char*)mHostMemoryBase + mHostMemorySize, mHostMemoryPoolEnd), ptrDiff(mHostMemoryPool, mHostMemoryBase), mHostMemoryPoolBlocked == nullptr ? 0ll : ptrDiff((char*)mHostMemoryBase + mHostMemorySize, mHostMemoryPoolBlocked)); printf("%16s", ""); PrintMemoryMax(); } diff --git a/GPU/GPUTracking/Base/GPUReconstruction.h b/GPU/GPUTracking/Base/GPUReconstruction.h index 33c9a9e0f9a94..26a4135a22610 100644 --- a/GPU/GPUTracking/Base/GPUReconstruction.h +++ b/GPU/GPUTracking/Base/GPUReconstruction.h @@ -201,8 +201,8 @@ class GPUReconstruction void FreeRegisteredMemory(short res); void ClearAllocatedMemory(bool clearOutputs = true); void ReturnVolatileDeviceMemory(); - void PushNonPersistentMemory(); - void PopNonPersistentMemory(RecoStep step); + void PushNonPersistentMemory(unsigned long tag); + void PopNonPersistentMemory(RecoStep step, unsigned long tag); void BlockStackedMemory(GPUReconstruction* rec); void UnblockStackedMemory(); void ResetRegisteredMemoryPointers(GPUProcessor* proc); @@ -385,7 +385,7 @@ class GPUReconstruction std::vector res; }; std::unordered_map mMemoryReuse1to1; - std::vector> mNonPersistentMemoryStack; + std::vector> mNonPersistentMemoryStack; std::vector mNonPersistentIndividualAllocations; std::unique_ptr mPipelineContext; diff --git a/GPU/GPUTracking/Global/GPUChainTracking.cxx b/GPU/GPUTracking/Global/GPUChainTracking.cxx index 23a38a2d25174..237489d57d02a 100644 --- a/GPU/GPUTracking/Global/GPUChainTracking.cxx +++ b/GPU/GPUTracking/Global/GPUChainTracking.cxx @@ -52,6 +52,7 @@ #include "TPCFastTransform.h" #include "utils/linux_helpers.h" +#include "utils/strtag.h" using namespace GPUCA_NAMESPACE::gpu; #include "GPUO2DataTypes.h" @@ -564,7 +565,7 @@ int GPUChainTracking::RunChain() return 1; } - mRec->PushNonPersistentMemory(); // 1st stack level for TPC tracking slice data + mRec->PushNonPersistentMemory(qStr2Tag("TPCSLCD1")); // 1st stack level for TPC tracking slice data if (runRecoStep(RecoStep::TPCSliceTracking, &GPUChainTracking::RunTPCTrackingSlices)) { return 1; } @@ -576,7 +577,7 @@ int GPUChainTracking::RunChain() if (runRecoStep(RecoStep::TPCMerging, &GPUChainTracking::RunTPCTrackingMerger, false)) { return 1; } - mRec->PopNonPersistentMemory(RecoStep::TPCSliceTracking); // Release 1st stack level, TPC slice data not needed after merger + mRec->PopNonPersistentMemory(RecoStep::TPCSliceTracking, qStr2Tag("TPCSLCD1")); // Release 1st stack level, TPC slice data not needed after merger if (mIOPtrs.clustersNative) { if (GetProcessingSettings().doublePipeline) { diff --git a/GPU/GPUTracking/Global/GPUChainTrackingClusterizer.cxx b/GPU/GPUTracking/Global/GPUChainTrackingClusterizer.cxx index 8552c1f1ccfd2..61b6d650f193d 100644 --- a/GPU/GPUTracking/Global/GPUChainTrackingClusterizer.cxx +++ b/GPU/GPUTracking/Global/GPUChainTrackingClusterizer.cxx @@ -34,6 +34,8 @@ #include "GPUO2FakeClasses.h" #endif +#include "utils/strtag.h" + #ifndef GPUCA_NO_VC #include #endif @@ -374,7 +376,7 @@ int GPUChainTracking::RunTPCClusterizer(bool synchronizeOutput) return ForwardTPCDigits(); } #ifdef GPUCA_TPC_GEOMETRY_O2 - mRec->PushNonPersistentMemory(); + mRec->PushNonPersistentMemory(qStr2Tag("TPCCLUST")); const auto& threadContext = GetThreadContext(); bool doGPU = GetRecoStepsGPU() & RecoStep::TPCClusterFinding; if (RunTPCClusterizer_prepare(mPipelineNotifyCtx && GetProcessingSettings().doublePipelineClusterizer)) { @@ -739,7 +741,7 @@ int GPUChainTracking::RunTPCClusterizer(bool synchronizeOutput) } } mRec->MemoryScalers()->nTPCHits = nClsTotal; - mRec->PopNonPersistentMemory(RecoStep::TPCClusterFinding); + mRec->PopNonPersistentMemory(RecoStep::TPCClusterFinding, qStr2Tag("TPCCLUST")); if (mPipelineNotifyCtx) { mRec->UnblockStackedMemory(); mPipelineNotifyCtx = nullptr; diff --git a/GPU/GPUTracking/Global/GPUChainTrackingCompression.cxx b/GPU/GPUTracking/Global/GPUChainTrackingCompression.cxx index 2bdc14b11ccf6..031758f71fadd 100644 --- a/GPU/GPUTracking/Global/GPUChainTrackingCompression.cxx +++ b/GPU/GPUTracking/Global/GPUChainTrackingCompression.cxx @@ -20,6 +20,7 @@ #include "GPUTPCCFChainContext.h" #include "TPCClusterDecompressor.h" #endif +#include "utils/strtag.h" using namespace GPUCA_NAMESPACE::gpu; using namespace o2::tpc; @@ -27,7 +28,7 @@ using namespace o2::tpc; int GPUChainTracking::RunTPCCompression() { #ifdef HAVE_O2HEADERS - mRec->PushNonPersistentMemory(); + mRec->PushNonPersistentMemory(qStr2Tag("TPCCOMPR")); RecoStep myStep = RecoStep::TPCCompression; bool doGPU = GetRecoStepsGPU() & RecoStep::TPCCompression; GPUTPCCompression& Compressor = processors()->tpcCompressor; @@ -205,7 +206,7 @@ int GPUChainTracking::RunTPCCompression() } else { ((GPUChainTracking*)GetNextChainInQueue())->mRec->BlockStackedMemory(mRec); } - mRec->PopNonPersistentMemory(RecoStep::TPCCompression); + mRec->PopNonPersistentMemory(RecoStep::TPCCompression, qStr2Tag("TPCCOMPR")); #endif return 0; } diff --git a/GPU/GPUTracking/Global/GPUChainTrackingMerger.cxx b/GPU/GPUTracking/Global/GPUChainTrackingMerger.cxx index 4fab18d788ecd..df3c1b593625a 100644 --- a/GPU/GPUTracking/Global/GPUChainTrackingMerger.cxx +++ b/GPU/GPUTracking/Global/GPUChainTrackingMerger.cxx @@ -14,6 +14,7 @@ #include "GPUChainTracking.h" #include "GPULogging.h" #include "GPUO2DataTypes.h" +#include "utils/strtag.h" #include using namespace GPUCA_NAMESPACE::gpu; @@ -95,7 +96,7 @@ int GPUChainTracking::RunTPCTrackingMerger(bool synchronizeOutput) TransferMemoryResourcesToGPU(RecoStep::NoRecoStep, &trk, 0); } } - mRec->PushNonPersistentMemory(); + mRec->PushNonPersistentMemory(qStr2Tag("TPCMERGE")); bool doGPU = GetRecoStepsGPU() & RecoStep::TPCMerging; bool doGPUall = doGPU && GetProcessingSettings().fullMergerOnGPU; GPUReconstruction::krnlDeviceType deviceType = doGPUall ? GPUReconstruction::krnlDeviceType::Auto : GPUReconstruction::krnlDeviceType::CPU; @@ -316,6 +317,6 @@ int GPUChainTracking::RunTPCTrackingMerger(bool synchronizeOutput) if (GetProcessingSettings().debugLevel >= 2) { GPUInfo("TPC Merger Finished (output clusters %d / input clusters %d)", Merger.NOutputTrackClusters(), Merger.NClusters()); } - mRec->PopNonPersistentMemory(RecoStep::TPCMerging); + mRec->PopNonPersistentMemory(RecoStep::TPCMerging, qStr2Tag("TPCMERGE")); return 0; } diff --git a/GPU/GPUTracking/Global/GPUChainTrackingSliceTracker.cxx b/GPU/GPUTracking/Global/GPUChainTrackingSliceTracker.cxx index d6dc1552bf620..45c28bb8433f6 100644 --- a/GPU/GPUTracking/Global/GPUChainTrackingSliceTracker.cxx +++ b/GPU/GPUTracking/Global/GPUChainTrackingSliceTracker.cxx @@ -16,6 +16,7 @@ #include "GPUO2DataTypes.h" #include "GPUMemorySizeScalers.h" #include "GPUTPCClusterData.h" +#include "utils/strtag.h" #include using namespace GPUCA_NAMESPACE::gpu; @@ -99,7 +100,7 @@ int GPUChainTracking::RunTPCTrackingSlices_internal() mRec->AllocateRegisteredMemory(processors()->tpcTrackers[iSlice].MemoryResSliceScratch()); mRec->AllocateRegisteredMemory(processors()->tpcTrackers[iSlice].MemoryResSliceInput()); } - mRec->PushNonPersistentMemory(); + mRec->PushNonPersistentMemory(qStr2Tag("TPCSLTRK")); for (unsigned int iSlice = 0; iSlice < NSLICES; iSlice++) { SetupGPUProcessor(&processors()->tpcTrackers[iSlice], true); // Now we allocate mRec->ResetRegisteredMemoryPointers(&processors()->tpcTrackers[iSlice]); // TODO: The above call breaks the GPU ptrs to already allocated memory. This fixes them. Should actually be cleaned up at the source. @@ -462,7 +463,7 @@ int GPUChainTracking::RunTPCTrackingSlices_internal() if (GetProcessingSettings().debugLevel >= 2) { GPUInfo("TPC Slice Tracker finished"); } - mRec->PopNonPersistentMemory(RecoStep::TPCSliceTracking); + mRec->PopNonPersistentMemory(RecoStep::TPCSliceTracking, qStr2Tag("TPCSLTRK")); return 0; } diff --git a/GPU/GPUTracking/Global/GPUChainTrackingTRD.cxx b/GPU/GPUTracking/Global/GPUChainTrackingTRD.cxx index 7c3dd03305421..63c39b0d69474 100644 --- a/GPU/GPUTracking/Global/GPUChainTrackingTRD.cxx +++ b/GPU/GPUTracking/Global/GPUChainTrackingTRD.cxx @@ -18,6 +18,7 @@ #include "GPUTRDTrackletLabels.h" #include "GPUTRDTrack.h" #include "GPUTRDTracker.h" +#include "utils/strtag.h" using namespace GPUCA_NAMESPACE::gpu; using namespace o2::trd; @@ -34,7 +35,7 @@ int GPUChainTracking::RunTRDTracking() return 0; } - mRec->PushNonPersistentMemory(); + mRec->PushNonPersistentMemory(qStr2Tag("TRDTRACK")); SetupGPUProcessor(&Tracker, true); for (unsigned int iTracklet = 0; iTracklet < mIOPtrs.nTRDTracklets; ++iTracklet) { @@ -62,7 +63,7 @@ int GPUChainTracking::RunTRDTracking() mIOPtrs.nTRDTracks = Tracker.NTracks(); mIOPtrs.trdTracks = Tracker.Tracks(); - mRec->PopNonPersistentMemory(RecoStep::TRDTracking); + mRec->PopNonPersistentMemory(RecoStep::TRDTracking, qStr2Tag("TRDTRACK")); return 0; } diff --git a/GPU/GPUTracking/Global/GPUChainTrackingTransformation.cxx b/GPU/GPUTracking/Global/GPUChainTrackingTransformation.cxx index b66c784041058..deb2843451e6a 100644 --- a/GPU/GPUTracking/Global/GPUChainTrackingTransformation.cxx +++ b/GPU/GPUTracking/Global/GPUChainTrackingTransformation.cxx @@ -25,6 +25,7 @@ #else #include "GPUO2FakeClasses.h" #endif +#include "utils/strtag.h" using namespace GPUCA_NAMESPACE::gpu; using namespace o2::tpc; @@ -32,7 +33,7 @@ using namespace o2::tpc; int GPUChainTracking::ConvertNativeToClusterData() { #ifdef HAVE_O2HEADERS - mRec->PushNonPersistentMemory(); + mRec->PushNonPersistentMemory(qStr2Tag("TPCTRANS")); const auto& threadContext = GetThreadContext(); bool doGPU = GetRecoStepsGPU() & RecoStep::TPCConversion; GPUTPCConvert& convert = processors()->tpcConverter; @@ -72,7 +73,7 @@ int GPUChainTracking::ConvertNativeToClusterData() mIOPtrs.nClusterData[i] = (i == NSLICES - 1 ? mIOPtrs.clustersNative->nClustersTotal : mIOPtrs.clustersNative->clusterOffset[i + 1][0]) - mIOPtrs.clustersNative->clusterOffset[i][0]; mIOPtrs.clusterData[i] = convert.mClusters + mIOPtrs.clustersNative->clusterOffset[i][0]; } - mRec->PopNonPersistentMemory(RecoStep::TPCConversion); + mRec->PopNonPersistentMemory(RecoStep::TPCConversion, qStr2Tag("TPCTRANS")); #endif return 0; } From dde3f22030fa8bb2e6aad27c583ed03f7a24be06 Mon Sep 17 00:00:00 2001 From: David Rohr Date: Wed, 3 Mar 2021 19:14:46 +0100 Subject: [PATCH 16/20] GPU: Reduce TPC Track merging memory footprint, can free scratch memory before final output is reformatted --- GPU/GPUTracking/Global/GPUChainTracking.cxx | 6 +++++- GPU/GPUTracking/Global/GPUChainTracking.h | 4 +++- .../Global/GPUChainTrackingMerger.cxx | 17 ++++++++++++++--- GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx | 12 ++++++++---- GPU/GPUTracking/Merger/GPUTPCGMMerger.h | 3 +++ 5 files changed, 33 insertions(+), 9 deletions(-) diff --git a/GPU/GPUTracking/Global/GPUChainTracking.cxx b/GPU/GPUTracking/Global/GPUChainTracking.cxx index 237489d57d02a..3e2149dd92877 100644 --- a/GPU/GPUTracking/Global/GPUChainTracking.cxx +++ b/GPU/GPUTracking/Global/GPUChainTracking.cxx @@ -566,6 +566,7 @@ int GPUChainTracking::RunChain() } mRec->PushNonPersistentMemory(qStr2Tag("TPCSLCD1")); // 1st stack level for TPC tracking slice data + mTPCSliceScratchOnStack = true; if (runRecoStep(RecoStep::TPCSliceTracking, &GPUChainTracking::RunTPCTrackingSlices)) { return 1; } @@ -577,7 +578,10 @@ int GPUChainTracking::RunChain() if (runRecoStep(RecoStep::TPCMerging, &GPUChainTracking::RunTPCTrackingMerger, false)) { return 1; } - mRec->PopNonPersistentMemory(RecoStep::TPCSliceTracking, qStr2Tag("TPCSLCD1")); // Release 1st stack level, TPC slice data not needed after merger + if (mTPCSliceScratchOnStack) { + mRec->PopNonPersistentMemory(RecoStep::TPCSliceTracking, qStr2Tag("TPCSLCD1")); // Release 1st stack level, TPC slice data not needed after merger + mTPCSliceScratchOnStack = false; + } if (mIOPtrs.clustersNative) { if (GetProcessingSettings().doublePipeline) { diff --git a/GPU/GPUTracking/Global/GPUChainTracking.h b/GPU/GPUTracking/Global/GPUChainTracking.h index ba9493d562888..a85052adfa425 100644 --- a/GPU/GPUTracking/Global/GPUChainTracking.h +++ b/GPU/GPUTracking/Global/GPUChainTracking.h @@ -250,11 +250,13 @@ class GPUChainTracking : public GPUChain, GPUReconstructionHelpers::helperDelega std::unique_ptr mMatLUTU; // Material Lookup Table std::unique_ptr mTRDGeometryU; // TRD Geometry + // Ptrs to internal buffers std::unique_ptr mClusterNativeAccess; - std::array mSubOutputControls = {nullptr}; + // (Ptrs to) configuration objects std::unique_ptr mCFContext; + bool mTPCSliceScratchOnStack = false; // Upper bounds for memory allocation unsigned int mMaxTPCHits = 0; diff --git a/GPU/GPUTracking/Global/GPUChainTrackingMerger.cxx b/GPU/GPUTracking/Global/GPUChainTrackingMerger.cxx index df3c1b593625a..5bbd34c40338e 100644 --- a/GPU/GPUTracking/Global/GPUChainTrackingMerger.cxx +++ b/GPU/GPUTracking/Global/GPUChainTrackingMerger.cxx @@ -259,12 +259,23 @@ int GPUChainTracking::RunTPCTrackingMerger(bool synchronizeOutput) TransferMemoryResourcesToHost(RecoStep::TPCMerging, &Merger, -1, true); } + mRec->ReturnVolatileDeviceMemory(); + mRec->PopNonPersistentMemory(RecoStep::TPCMerging, qStr2Tag("TPCMERGE")); + #ifdef GPUCA_TPC_GEOMETRY_O2 if (GetProcessingSettings().createO2Output) { - mRec->ReturnVolatileDeviceMemory(); + if (mTPCSliceScratchOnStack) { + mRec->PopNonPersistentMemory(RecoStep::TPCSliceTracking, qStr2Tag("TPCSLCD1")); // Return the slice data memory early + mTPCSliceScratchOnStack = false; + } + + mRec->PushNonPersistentMemory(qStr2Tag("TPCMERG2")); + AllocateRegisteredMemory(Merger.MemoryResOutputO2Scratch()); + WriteToConstantMemory(RecoStep::TPCMerging, (char*)&processors()->tpcMerger - (char*)processors(), &MergerShadow, sizeof(MergerShadow), 0); runKernel(GetGridAuto(0, deviceType), krnlRunRangeNone, krnlEventNone); TransferMemoryResourceLinkToHost(RecoStep::TPCMerging, Merger.MemoryResMemory(), 0); SynchronizeStream(0); + AllocateRegisteredMemory(Merger.MemoryResOutputO2(), mSubOutputControls[GPUTrackingOutputs::getIndex(&GPUTrackingOutputs::tpcTracksO2)]); AllocateRegisteredMemory(Merger.MemoryResOutputO2Clus(), mSubOutputControls[GPUTrackingOutputs::getIndex(&GPUTrackingOutputs::tpcTracksO2ClusRefs)]); WriteToConstantMemory(RecoStep::TPCMerging, (char*)&processors()->tpcMerger - (char*)processors(), &MergerShadow, sizeof(MergerShadow), 0); @@ -281,13 +292,14 @@ int GPUChainTracking::RunTPCTrackingMerger(bool synchronizeOutput) TransferMemoryResourceLinkToHost(RecoStep::TPCMerging, Merger.MemoryResOutputO2Clus(), mRec->NStreams() - 2, nullptr, &mEvents->single); ReleaseEvent(&mEvents->single); } + mRec->ReturnVolatileDeviceMemory(); + mRec->PopNonPersistentMemory(RecoStep::TPCMerging, qStr2Tag("TPCMERG2")); } #endif if (synchronizeOutput) { SynchronizeStream(mRec->NStreams() - 2); } - mRec->ReturnVolatileDeviceMemory(); mIOPtrs.mergedTracks = Merger.OutputTracks(); mIOPtrs.nMergedTracks = Merger.NOutputTracks(); mIOPtrs.mergedTrackHits = Merger.Clusters(); @@ -317,6 +329,5 @@ int GPUChainTracking::RunTPCTrackingMerger(bool synchronizeOutput) if (GetProcessingSettings().debugLevel >= 2) { GPUInfo("TPC Merger Finished (output clusters %d / input clusters %d)", Merger.NOutputTrackClusters(), Merger.NClusters()); } - mRec->PopNonPersistentMemory(RecoStep::TPCMerging, qStr2Tag("TPCMERGE")); return 0; } diff --git a/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx b/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx index a197526f2e95f..3f32cefeadb22 100644 --- a/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx +++ b/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx @@ -264,10 +264,6 @@ void* GPUTPCGMMerger::SetPointersMerger(void* mem) computePointerWithAlignment(mem, mTrackCCRoots, mNMaxSliceTracks); void* memMax = mem; mem = memBase; - computePointerWithAlignment(mem, mTrackSortO2, mNMaxTracks); // GPUTPCGMO2Output::prepare - GPUTPCGMO2Output::sort - GPUTPCGMO2Output::output - computePointerWithAlignment(mem, mClusRefTmp, mNMaxTracks); - memMax = (void*)std::max((size_t)mem, (size_t)memMax); // BUG: Casts needed to workaround GCC miscompilation - mem = memBase; computePointerWithAlignment(mem, mTrackIDs, mNMaxTracks); // UnpackResetIds - RefitSliceTracks - UnpackSliceGlobal memMax = (void*)std::max((size_t)mem, (size_t)memMax); mem = memBase; @@ -349,6 +345,13 @@ void* GPUTPCGMMerger::SetPointersOutputO2MC(void* mem) return mem; } +void* GPUTPCGMMerger::SetPointersOutputO2Scratch(void* mem) +{ + computePointerWithAlignment(mem, mTrackSortO2, mNMaxTracks); + computePointerWithAlignment(mem, mClusRefTmp, mNMaxTracks); + return mem; +} + void GPUTPCGMMerger::RegisterMemoryAllocation() { AllocateAndInitializeLate(); @@ -357,6 +360,7 @@ void GPUTPCGMMerger::RegisterMemoryAllocation() mMemoryResOutput = mRec->RegisterMemoryAllocation(this, &GPUTPCGMMerger::SetPointersOutput, (mRec->GetProcessingSettings().fullMergerOnGPU ? (mRec->GetProcessingSettings().createO2Output > 1 ? GPUMemoryResource::MEMORY_SCRATCH : GPUMemoryResource::MEMORY_OUTPUT) : GPUMemoryResource::MEMORY_INOUT) | GPUMemoryResource::MEMORY_CUSTOM, "TPCMergerOutput"); mMemoryResOutputState = mRec->RegisterMemoryAllocation(this, &GPUTPCGMMerger::SetPointersOutputState, (mRec->GetProcessingSettings().fullMergerOnGPU ? GPUMemoryResource::MEMORY_OUTPUT : GPUMemoryResource::MEMORY_HOST) | GPUMemoryResource::MEMORY_CUSTOM, "TPCMergerOutputState"); if (mRec->GetProcessingSettings().createO2Output) { + mMemoryResOutputO2Scratch = mRec->RegisterMemoryAllocation(this, &GPUTPCGMMerger::SetPointersOutputO2Scratch, GPUMemoryResource::MEMORY_SCRATCH | GPUMemoryResource::MEMORY_STACK | GPUMemoryResource::MEMORY_CUSTOM, "TPCMergerOutputO2Scratch"); mMemoryResOutputO2 = mRec->RegisterMemoryAllocation(this, &GPUTPCGMMerger::SetPointersOutputO2, GPUMemoryResource::MEMORY_OUTPUT | GPUMemoryResource::MEMORY_CUSTOM, "TPCMergerOutputO2"); mMemoryResOutputO2Clus = mRec->RegisterMemoryAllocation(this, &GPUTPCGMMerger::SetPointersOutputO2Clus, GPUMemoryResource::MEMORY_OUTPUT | GPUMemoryResource::MEMORY_CUSTOM, "TPCMergerOutputO2Clus"); if (mRec->GetProcessingSettings().runMC) { diff --git a/GPU/GPUTracking/Merger/GPUTPCGMMerger.h b/GPU/GPUTracking/Merger/GPUTPCGMMerger.h index 44c3599387176..500419741716d 100644 --- a/GPU/GPUTracking/Merger/GPUTPCGMMerger.h +++ b/GPU/GPUTracking/Merger/GPUTPCGMMerger.h @@ -96,6 +96,7 @@ class GPUTPCGMMerger : public GPUProcessor void* SetPointersOutputO2(void* mem); void* SetPointersOutputO2Clus(void* mem); void* SetPointersOutputO2MC(void* mem); + void* SetPointersOutputO2Scratch(void* mem); void* SetPointersOutputState(void* mem); void* SetPointersMemory(void* mem); @@ -140,6 +141,7 @@ class GPUTPCGMMerger : public GPUProcessor GPUd() unsigned short MemoryResOutputO2() const { return mMemoryResOutputO2; } GPUd() unsigned short MemoryResOutputO2Clus() const { return mMemoryResOutputO2Clus; } GPUd() unsigned short MemoryResOutputO2MC() const { return mMemoryResOutputO2MC; } + GPUd() unsigned short MemoryResOutputO2Scratch() const { return mMemoryResOutputO2Scratch; } GPUd() int RefitSliceTrack(GPUTPCGMSliceTrack& sliceTrack, const GPUTPCTrack* inTrack, float alpha, int slice); GPUd() void SetTrackClusterZT(GPUTPCGMSliceTrack& track, int iSlice, const GPUTPCTrack* sliceTr); @@ -236,6 +238,7 @@ class GPUTPCGMMerger : public GPUProcessor unsigned short mMemoryResOutputO2; unsigned short mMemoryResOutputO2Clus; unsigned short mMemoryResOutputO2MC; + unsigned short mMemoryResOutputO2Scratch; int mNClusters; // Total number of incoming clusters (from slice tracks) GPUTPCGMMergedTrack* mOutputTracks; //* array of output merged tracks From c28f728e93b756f6da2b762b69c4b6168b945ddf Mon Sep 17 00:00:00 2001 From: David Rohr Date: Thu, 4 Mar 2021 09:45:15 +0100 Subject: [PATCH 17/20] GPU: Add possibility to user smaller unattached compressed hits buffer in TPC compression than total number of clusters --- .../include/DataFormatsTPC/ClusterNative.h | 2 +- .../DataCompression/GPUTPCCompression.cxx | 7 +- .../DataCompression/GPUTPCCompression.h | 3 +- .../GPUTPCCompressionKernels.cxx | 64 +++++++++++-------- .../GPUTPCCompressionKernels.h | 2 +- .../DataTypes/GPUMemorySizeScalers.h | 1 + .../Global/GPUChainTrackingCompression.cxx | 23 +++---- .../GPUChainTrackingDebugAndProfiling.cxx | 5 ++ GPU/GPUTracking/Global/GPUErrorCodes.h | 1 + 9 files changed, 60 insertions(+), 48 deletions(-) diff --git a/DataFormats/Detectors/TPC/include/DataFormatsTPC/ClusterNative.h b/DataFormats/Detectors/TPC/include/DataFormatsTPC/ClusterNative.h index 70124c72d62c9..beacd5b4f2909 100644 --- a/DataFormats/Detectors/TPC/include/DataFormatsTPC/ClusterNative.h +++ b/DataFormats/Detectors/TPC/include/DataFormatsTPC/ClusterNative.h @@ -164,7 +164,7 @@ struct ClusterNativeAccess { unsigned int nClusters[constants::MAXSECTOR][constants::MAXGLOBALPADROW]; unsigned int nClustersSector[constants::MAXSECTOR]; unsigned int clusterOffset[constants::MAXSECTOR][constants::MAXGLOBALPADROW]; - unsigned int nClustersTotal; + unsigned int nClustersTotal; // Must be directly after clusterOffsets, --> =clusterOffset[nRows * nSectors]! void setOffsetPtrs(); diff --git a/GPU/GPUTracking/DataCompression/GPUTPCCompression.cxx b/GPU/GPUTracking/DataCompression/GPUTPCCompression.cxx index 02f141fe7baee..f87465f8d967f 100644 --- a/GPU/GPUTracking/DataCompression/GPUTPCCompression.cxx +++ b/GPU/GPUTracking/DataCompression/GPUTPCCompression.cxx @@ -14,6 +14,7 @@ #include "GPUTPCCompression.h" #include "GPUReconstruction.h" #include "GPUO2DataTypes.h" +#include "GPUMemorySizeScalers.h" using namespace GPUCA_NAMESPACE::gpu; @@ -39,7 +40,7 @@ void* GPUTPCCompression::SetPointersScratch(void* mem) computePointerWithAlignment(mem, mAttachedClusterFirstIndex, mMaxTracks); } if (mRec->GetProcessingSettings().tpcCompressionGatherMode != 1) { - SetPointersCompressedClusters(mem, mPtrs, mMaxTrackClusters, mMaxTracks, mMaxClusters, false); + SetPointersCompressedClusters(mem, mPtrs, mMaxTrackClusters, mMaxTracks, mMaxClustersInCache, false); } return mem; } @@ -48,7 +49,7 @@ void* GPUTPCCompression::SetPointersOutput(void* mem) { computePointerWithAlignment(mem, mAttachedClusterFirstIndex, mMaxTrackClusters); if (mRec->GetProcessingSettings().tpcCompressionGatherMode == 1) { - SetPointersCompressedClusters(mem, mPtrs, mMaxTrackClusters, mMaxTracks, mMaxClusters, false); + SetPointersCompressedClusters(mem, mPtrs, mMaxTrackClusters, mMaxTracks, mMaxClustersInCache, false); } return mem; } @@ -115,6 +116,8 @@ void GPUTPCCompression::RegisterMemoryAllocation() void GPUTPCCompression::SetMaxData(const GPUTrackingInOutPointers& io) { mMaxClusters = io.clustersNative->nClustersTotal; + mMaxClusterFactorBase1024 = mMaxClusters > 100000000 ? mRec->MemoryScalers()->tpcCompressedUnattachedHitsBase1024[mRec->GetParam().rec.tpcRejectionMode] : 1024; + mMaxClustersInCache = mMaxClusters * mMaxClusterFactorBase1024 / 1024; mMaxTrackClusters = mRec->GetConstantMem().tpcMerger.NOutputTrackClusters(); mMaxTracks = mRec->GetConstantMem().tpcMerger.NOutputTracks(); if (mMaxClusters % 16) { diff --git a/GPU/GPUTracking/DataCompression/GPUTPCCompression.h b/GPU/GPUTracking/DataCompression/GPUTPCCompression.h index b0f449bd81edb..73eab9db2d325 100644 --- a/GPU/GPUTracking/DataCompression/GPUTPCCompression.h +++ b/GPU/GPUTracking/DataCompression/GPUTPCCompression.h @@ -90,7 +90,8 @@ class GPUTPCCompression : public GPUProcessor unsigned int mMaxTracks = 0; unsigned int mMaxClusters = 0; unsigned int mMaxTrackClusters = 0; - unsigned int mNMaxClusterSliceRow = 0; + unsigned int mMaxClustersInCache = 0; + size_t mMaxClusterFactorBase1024 = 0; template void SetPointersCompressedClusters(void*& mem, T& c, unsigned int nClA, unsigned int nTr, unsigned int nClU, bool reducedClA); diff --git a/GPU/GPUTracking/DataCompression/GPUTPCCompressionKernels.cxx b/GPU/GPUTracking/DataCompression/GPUTPCCompressionKernels.cxx index a00b1375416fd..a29c38a53bcad 100644 --- a/GPU/GPUTracking/DataCompression/GPUTPCCompressionKernels.cxx +++ b/GPU/GPUTracking/DataCompression/GPUTPCCompressionKernels.cxx @@ -187,6 +187,8 @@ GPUdii() void GPUTPCCompressionKernels::ThreadclusterOffset[iSlice][iRow]; + const int idOffsetOut = clusters->clusterOffset[iSlice][iRow] * compressor.mMaxClusterFactorBase1024 / 1024; + const int idOffsetOutMax = clusters->clusterOffset[iSlice][iRow + 1] * compressor.mMaxClusterFactorBase1024 / 1024; if (iThread == nThreads - 1) { smem.nCount = 0; } @@ -246,7 +248,7 @@ GPUdii() void GPUTPCCompressionKernels::Thread(clusters->clusters[iSlice][iRow])); @@ -261,7 +263,12 @@ GPUdii() void GPUTPCCompressionKernels::Thread= idOffsetOutMax) { + compressor.raiseError(GPUErrors::ERROR_COMPRESSION_ROW_HIT_OVERFLOW, outidx, idOffsetOutMax); + count = 0; + break; + } const ClusterNative& GPUrestrict() orgCl = clusters->clusters[iSlice][iRow][sortBuffer[j]]; unsigned int lastTime = 0; unsigned int lastPad = 0; @@ -446,7 +453,7 @@ GPUdi() void GPUTPCCompressionGatherKernels::compressorMemcpyBasic(T* GPUrestric } template -GPUdi() void GPUTPCCompressionGatherKernels::compressorMemcpyBuffered(V* buf, T* GPUrestrict() dst, const T* GPUrestrict() src, const S* GPUrestrict() nums, const unsigned int* GPUrestrict() srcOffsets, unsigned int nTracks, int nLanes, int iLane, int diff) +GPUdi() void GPUTPCCompressionGatherKernels::compressorMemcpyBuffered(V* buf, T* GPUrestrict() dst, const T* GPUrestrict() src, const S* GPUrestrict() nums, const unsigned int* GPUrestrict() srcOffsets, unsigned int nEntries, int nLanes, int iLane, int diff, size_t scaleBase1024) { int shmPos = 0; unsigned int dstOffset = 0; @@ -456,9 +463,9 @@ GPUdi() void GPUTPCCompressionGatherKernels::compressorMemcpyBuffered(V* buf, T* CONSTEXPR int bufSize = GPUCA_WARP_SIZE; CONSTEXPR int bufTSize = bufSize * sizeof(V) / sizeof(T); - for (unsigned int i = 0; i < nTracks; i++) { + for (unsigned int i = 0; i < nEntries; i++) { unsigned int srcPos = 0; - unsigned int srcOffset = srcOffsets[i] + diff; + unsigned int srcOffset = (srcOffsets[i] * scaleBase1024 / 1024) + diff; unsigned int srcSize = nums[i] - diff; if (dstAligned == nullptr) { @@ -565,13 +572,14 @@ GPUdii() void GPUTPCCompressionGatherKernels::ThreadqTotU + rowsOffset, compressor.mPtrs.qTotU + clusters->clusterOffset[i][j], nClusters, nLanes, iLane); - compressorMemcpy(compressor.mOutput->qMaxU + rowsOffset, compressor.mPtrs.qMaxU + clusters->clusterOffset[i][j], nClusters, nLanes, iLane); - compressorMemcpy(compressor.mOutput->flagsU + rowsOffset, compressor.mPtrs.flagsU + clusters->clusterOffset[i][j], nClusters, nLanes, iLane); - compressorMemcpy(compressor.mOutput->padDiffU + rowsOffset, compressor.mPtrs.padDiffU + clusters->clusterOffset[i][j], nClusters, nLanes, iLane); - compressorMemcpy(compressor.mOutput->timeDiffU + rowsOffset, compressor.mPtrs.timeDiffU + clusters->clusterOffset[i][j], nClusters, nLanes, iLane); - compressorMemcpy(compressor.mOutput->sigmaPadU + rowsOffset, compressor.mPtrs.sigmaPadU + clusters->clusterOffset[i][j], nClusters, nLanes, iLane); - compressorMemcpy(compressor.mOutput->sigmaTimeU + rowsOffset, compressor.mPtrs.sigmaTimeU + clusters->clusterOffset[i][j], nClusters, nLanes, iLane); + unsigned int clusterOffsetInCache = clusters->clusterOffset[i][j] * compressor.mMaxClusterFactorBase1024 / 1024; + compressorMemcpy(compressor.mOutput->qTotU + rowsOffset, compressor.mPtrs.qTotU + clusterOffsetInCache, nClusters, nLanes, iLane); + compressorMemcpy(compressor.mOutput->qMaxU + rowsOffset, compressor.mPtrs.qMaxU + clusterOffsetInCache, nClusters, nLanes, iLane); + compressorMemcpy(compressor.mOutput->flagsU + rowsOffset, compressor.mPtrs.flagsU + clusterOffsetInCache, nClusters, nLanes, iLane); + compressorMemcpy(compressor.mOutput->padDiffU + rowsOffset, compressor.mPtrs.padDiffU + clusterOffsetInCache, nClusters, nLanes, iLane); + compressorMemcpy(compressor.mOutput->timeDiffU + rowsOffset, compressor.mPtrs.timeDiffU + clusterOffsetInCache, nClusters, nLanes, iLane); + compressorMemcpy(compressor.mOutput->sigmaPadU + rowsOffset, compressor.mPtrs.sigmaPadU + clusterOffsetInCache, nClusters, nLanes, iLane); + compressorMemcpy(compressor.mOutput->sigmaTimeU + rowsOffset, compressor.mPtrs.sigmaTimeU + clusterOffsetInCache, nClusters, nLanes, iLane); rowsOffset += nClusters; } } @@ -676,18 +684,18 @@ GPUdii() void GPUTPCCompressionGatherKernels::gatherBuffered(int nBlocks, int nT compressorMemcpyBasic(output->padA, input.padA, compressor.mMemory->nStoredTracks, nThreads, iThread); } - const unsigned int* clusterOffsets = reinterpret_cast(clusters->clusterOffset) + rowStart; + const unsigned int* clusterOffsets = &clusters->clusterOffset[0][0] + rowStart; const unsigned int* nSliceRowClusters = input.nSliceRowClusters + rowStart; auto* buf = smem.getBuffer(iWarp); - compressorMemcpyBuffered(buf, output->qTotU + rowsOffset, input.qTotU, nSliceRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0); - compressorMemcpyBuffered(buf, output->qMaxU + rowsOffset, input.qMaxU, nSliceRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0); - compressorMemcpyBuffered(buf, output->flagsU + rowsOffset, input.flagsU, nSliceRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0); - compressorMemcpyBuffered(buf, output->padDiffU + rowsOffset, input.padDiffU, nSliceRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0); - compressorMemcpyBuffered(buf, output->timeDiffU + rowsOffset, input.timeDiffU, nSliceRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0); - compressorMemcpyBuffered(buf, output->sigmaPadU + rowsOffset, input.sigmaPadU, nSliceRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0); - compressorMemcpyBuffered(buf, output->sigmaTimeU + rowsOffset, input.sigmaTimeU, nSliceRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0); + compressorMemcpyBuffered(buf, output->qTotU + rowsOffset, input.qTotU, nSliceRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0, compressor.mMaxClusterFactorBase1024); + compressorMemcpyBuffered(buf, output->qMaxU + rowsOffset, input.qMaxU, nSliceRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0, compressor.mMaxClusterFactorBase1024); + compressorMemcpyBuffered(buf, output->flagsU + rowsOffset, input.flagsU, nSliceRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0, compressor.mMaxClusterFactorBase1024); + compressorMemcpyBuffered(buf, output->padDiffU + rowsOffset, input.padDiffU, nSliceRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0, compressor.mMaxClusterFactorBase1024); + compressorMemcpyBuffered(buf, output->timeDiffU + rowsOffset, input.timeDiffU, nSliceRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0, compressor.mMaxClusterFactorBase1024); + compressorMemcpyBuffered(buf, output->sigmaPadU + rowsOffset, input.sigmaPadU, nSliceRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0, compressor.mMaxClusterFactorBase1024); + compressorMemcpyBuffered(buf, output->sigmaTimeU + rowsOffset, input.sigmaTimeU, nSliceRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0, compressor.mMaxClusterFactorBase1024); const unsigned short* nTrackClustersPtr = input.nTrackClusters + trackStart; const unsigned int* aClsFstIdx = compressor.mAttachedClusterFirstIndex + trackStart; @@ -742,16 +750,16 @@ GPUdii() void GPUTPCCompressionGatherKernels::gatherMulti(int nBlocks, int nThre rowsPerWarp = rowEnd - rowStart; const unsigned int rowsOffset = calculateWarpOffsets(smem, input.nSliceRowClusters, rowStart, rowEnd, nWarps, iWarp, nLanes, iLane); - const unsigned int* clusterOffsets = reinterpret_cast(clusters->clusterOffset) + rowStart; + const unsigned int* clusterOffsets = &clusters->clusterOffset[0][0] + rowStart; const unsigned int* nSliceRowClusters = input.nSliceRowClusters + rowStart; - compressorMemcpyBuffered(buf, output->qTotU + rowsOffset, input.qTotU, nSliceRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0); - compressorMemcpyBuffered(buf, output->qMaxU + rowsOffset, input.qMaxU, nSliceRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0); - compressorMemcpyBuffered(buf, output->flagsU + rowsOffset, input.flagsU, nSliceRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0); - compressorMemcpyBuffered(buf, output->padDiffU + rowsOffset, input.padDiffU, nSliceRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0); - compressorMemcpyBuffered(buf, output->timeDiffU + rowsOffset, input.timeDiffU, nSliceRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0); - compressorMemcpyBuffered(buf, output->sigmaPadU + rowsOffset, input.sigmaPadU, nSliceRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0); - compressorMemcpyBuffered(buf, output->sigmaTimeU + rowsOffset, input.sigmaTimeU, nSliceRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0); + compressorMemcpyBuffered(buf, output->qTotU + rowsOffset, input.qTotU, nSliceRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0, compressor.mMaxClusterFactorBase1024); + compressorMemcpyBuffered(buf, output->qMaxU + rowsOffset, input.qMaxU, nSliceRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0, compressor.mMaxClusterFactorBase1024); + compressorMemcpyBuffered(buf, output->flagsU + rowsOffset, input.flagsU, nSliceRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0, compressor.mMaxClusterFactorBase1024); + compressorMemcpyBuffered(buf, output->padDiffU + rowsOffset, input.padDiffU, nSliceRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0, compressor.mMaxClusterFactorBase1024); + compressorMemcpyBuffered(buf, output->timeDiffU + rowsOffset, input.timeDiffU, nSliceRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0, compressor.mMaxClusterFactorBase1024); + compressorMemcpyBuffered(buf, output->sigmaPadU + rowsOffset, input.sigmaPadU, nSliceRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0, compressor.mMaxClusterFactorBase1024); + compressorMemcpyBuffered(buf, output->sigmaTimeU + rowsOffset, input.sigmaTimeU, nSliceRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0, compressor.mMaxClusterFactorBase1024); } else { const unsigned int nGlobalWarps = nWarps * (nBlocks - 1) / 2; const unsigned int iGlobalWarp = nWarps * (iBlock / 2 - 1) + iWarp; diff --git a/GPU/GPUTracking/DataCompression/GPUTPCCompressionKernels.h b/GPU/GPUTracking/DataCompression/GPUTPCCompressionKernels.h index c3fe9ce1e25ff..d7d14fc5d2ae6 100644 --- a/GPU/GPUTracking/DataCompression/GPUTPCCompressionKernels.h +++ b/GPU/GPUTracking/DataCompression/GPUTPCCompressionKernels.h @@ -112,7 +112,7 @@ class GPUTPCCompressionGatherKernels : public GPUKernelTemplate GPUdi() static void compressorMemcpyBasic(T* dst, const T* src, unsigned int size, int nThreads, int iThread, int nBlocks = 1, int iBlock = 0); template - GPUdi() static void compressorMemcpyBuffered(V* buf, T* dst, const T* src, const S* nums, const unsigned int* srcOffets, unsigned int nTracks, int nLanes, int iLane, int diff = 0); + GPUdi() static void compressorMemcpyBuffered(V* buf, T* dst, const T* src, const S* nums, const unsigned int* srcOffets, unsigned int nEntries, int nLanes, int iLane, int diff = 0, size_t scaleBase1024 = 1024); template GPUdi() static unsigned int calculateWarpOffsets(GPUSharedMemory& smem, T* nums, unsigned int start, unsigned int end, int nWarps, int iWarp, int nLanes, int iLane); diff --git a/GPU/GPUTracking/DataTypes/GPUMemorySizeScalers.h b/GPU/GPUTracking/DataTypes/GPUMemorySizeScalers.h index ed0aab483943b..542c5c35fc181 100644 --- a/GPU/GPUTracking/DataTypes/GPUMemorySizeScalers.h +++ b/GPU/GPUTracking/DataTypes/GPUMemorySizeScalers.h @@ -43,6 +43,7 @@ struct GPUMemorySizeScalers { double tpcSectorTrackHitsPerHit = 0.8f; double tpcMergedTrackPerSliceTrack = 0.9; double tpcMergedTrackHitPerSliceHit = 1.1; + size_t tpcCompressedUnattachedHitsBase1024[3] = {900, 900, 500}; // No ratio, but integer fraction of 1024 for exact computation // Upper limits size_t tpcMaxPeaks = 1000000000; diff --git a/GPU/GPUTracking/Global/GPUChainTrackingCompression.cxx b/GPU/GPUTracking/Global/GPUChainTrackingCompression.cxx index 031758f71fadd..fc72c20c22883 100644 --- a/GPU/GPUTracking/Global/GPUChainTrackingCompression.cxx +++ b/GPU/GPUTracking/Global/GPUChainTrackingCompression.cxx @@ -37,14 +37,6 @@ int GPUChainTracking::RunTPCCompression() if (mPipelineFinalizationCtx && GetProcessingSettings().doublePipelineClusterizer) { RecordMarker(&mEvents->single, 0); } - Compressor.mNMaxClusterSliceRow = 0; - for (unsigned int i = 0; i < NSLICES; i++) { - for (unsigned int j = 0; j < GPUCA_ROW_COUNT; j++) { - if (mIOPtrs.clustersNative->nClusters[i][j] > Compressor.mNMaxClusterSliceRow) { - Compressor.mNMaxClusterSliceRow = mIOPtrs.clustersNative->nClusters[i][j]; - } - } - } if (ProcessingSettings().tpcCompressionGatherMode == 3) { mRec->AllocateVolatileDeviceMemory(0); // make future device memory allocation volatile @@ -160,13 +152,14 @@ int GPUChainTracking::RunTPCCompression() unsigned int offset = 0; for (unsigned int i = 0; i < NSLICES; i++) { for (unsigned int j = 0; j < GPUCA_ROW_COUNT; j++) { - GPUMemCpyAlways(myStep, O->qTotU + offset, P->qTotU + mIOPtrs.clustersNative->clusterOffset[i][j], O->nSliceRowClusters[i * GPUCA_ROW_COUNT + j] * sizeof(O->qTotU[0]), outputStream, direction); - GPUMemCpyAlways(myStep, O->qMaxU + offset, P->qMaxU + mIOPtrs.clustersNative->clusterOffset[i][j], O->nSliceRowClusters[i * GPUCA_ROW_COUNT + j] * sizeof(O->qMaxU[0]), outputStream, direction); - GPUMemCpyAlways(myStep, O->flagsU + offset, P->flagsU + mIOPtrs.clustersNative->clusterOffset[i][j], O->nSliceRowClusters[i * GPUCA_ROW_COUNT + j] * sizeof(O->flagsU[0]), outputStream, direction); - GPUMemCpyAlways(myStep, O->padDiffU + offset, P->padDiffU + mIOPtrs.clustersNative->clusterOffset[i][j], O->nSliceRowClusters[i * GPUCA_ROW_COUNT + j] * sizeof(O->padDiffU[0]), outputStream, direction); - GPUMemCpyAlways(myStep, O->timeDiffU + offset, P->timeDiffU + mIOPtrs.clustersNative->clusterOffset[i][j], O->nSliceRowClusters[i * GPUCA_ROW_COUNT + j] * sizeof(O->timeDiffU[0]), outputStream, direction); - GPUMemCpyAlways(myStep, O->sigmaPadU + offset, P->sigmaPadU + mIOPtrs.clustersNative->clusterOffset[i][j], O->nSliceRowClusters[i * GPUCA_ROW_COUNT + j] * sizeof(O->sigmaPadU[0]), outputStream, direction); - GPUMemCpyAlways(myStep, O->sigmaTimeU + offset, P->sigmaTimeU + mIOPtrs.clustersNative->clusterOffset[i][j], O->nSliceRowClusters[i * GPUCA_ROW_COUNT + j] * sizeof(O->sigmaTimeU[0]), outputStream, direction); + unsigned int srcOffset = mIOPtrs.clustersNative->clusterOffset[i][j] * Compressor.mMaxClusterFactorBase1024 / 1024; + GPUMemCpyAlways(myStep, O->qTotU + offset, P->qTotU + srcOffset, O->nSliceRowClusters[i * GPUCA_ROW_COUNT + j] * sizeof(O->qTotU[0]), outputStream, direction); + GPUMemCpyAlways(myStep, O->qMaxU + offset, P->qMaxU + srcOffset, O->nSliceRowClusters[i * GPUCA_ROW_COUNT + j] * sizeof(O->qMaxU[0]), outputStream, direction); + GPUMemCpyAlways(myStep, O->flagsU + offset, P->flagsU + srcOffset, O->nSliceRowClusters[i * GPUCA_ROW_COUNT + j] * sizeof(O->flagsU[0]), outputStream, direction); + GPUMemCpyAlways(myStep, O->padDiffU + offset, P->padDiffU + srcOffset, O->nSliceRowClusters[i * GPUCA_ROW_COUNT + j] * sizeof(O->padDiffU[0]), outputStream, direction); + GPUMemCpyAlways(myStep, O->timeDiffU + offset, P->timeDiffU + srcOffset, O->nSliceRowClusters[i * GPUCA_ROW_COUNT + j] * sizeof(O->timeDiffU[0]), outputStream, direction); + GPUMemCpyAlways(myStep, O->sigmaPadU + offset, P->sigmaPadU + srcOffset, O->nSliceRowClusters[i * GPUCA_ROW_COUNT + j] * sizeof(O->sigmaPadU[0]), outputStream, direction); + GPUMemCpyAlways(myStep, O->sigmaTimeU + offset, P->sigmaTimeU + srcOffset, O->nSliceRowClusters[i * GPUCA_ROW_COUNT + j] * sizeof(O->sigmaTimeU[0]), outputStream, direction); offset += O->nSliceRowClusters[i * GPUCA_ROW_COUNT + j]; } } diff --git a/GPU/GPUTracking/Global/GPUChainTrackingDebugAndProfiling.cxx b/GPU/GPUTracking/Global/GPUChainTrackingDebugAndProfiling.cxx index 1083f34b7ebed..b13b11111aa00 100644 --- a/GPU/GPUTracking/Global/GPUChainTrackingDebugAndProfiling.cxx +++ b/GPU/GPUTracking/Global/GPUChainTrackingDebugAndProfiling.cxx @@ -149,11 +149,16 @@ void GPUChainTracking::PrintMemoryStatistics() addToMap("TPC Clusterer Clusters", usageMap, mRec->MemoryScalers()->nTPCHits, mRec->MemoryScalers()->NTPCClusters(mRec->MemoryScalers()->nTPCdigits)); addToMap("TPC Tracks", usageMap, processors()->tpcMerger.NOutputTracks(), processors()->tpcMerger.NMaxTracks()); addToMap("TPC TrackHits", usageMap, processors()->tpcMerger.NOutputTrackClusters(), processors()->tpcMerger.NMaxOutputTrackClusters()); + if (mRec->GetProcessingSettings().createO2Output) { addToMap("TPC O2 Tracks", usageMap, processors()->tpcMerger.NOutputTracksTPCO2(), processors()->tpcMerger.NOutputTracksTPCO2()); addToMap("TPC O2 ClusRefs", usageMap, processors()->tpcMerger.NOutputClusRefsTPCO2(), processors()->tpcMerger.NOutputClusRefsTPCO2()); } + addToMap("TPC ComprCache HitsAttached", usageMap, processors()->tpcCompressor.mOutput->nAttachedClusters, processors()->tpcCompressor.mMaxTrackClusters); + addToMap("TPC ComprCache HitsUnattached", usageMap, processors()->tpcCompressor.mOutput->nUnattachedClusters, processors()->tpcCompressor.mMaxClustersInCache); + addToMap("TPC ComprCache Tracks", usageMap, processors()->tpcCompressor.mOutput->nTracks, processors()->tpcCompressor.mMaxTracks); + for (auto& elem : usageMap) { printf("Mem Usage %-30s : %'14lu / %'14lu (%3.0f%% / %3.0f%% / count %3u / max %'14lu)\n", elem.first.c_str(), elem.second.nSum, elem.second.nBoundSum, 100. * elem.second.nSum / std::max(1lu, elem.second.nBoundSum), 100. * elem.second.maxUse, elem.second.count, elem.second.nMax); } diff --git a/GPU/GPUTracking/Global/GPUErrorCodes.h b/GPU/GPUTracking/Global/GPUErrorCodes.h index 07a616927e0ef..f6ae659e144d1 100644 --- a/GPU/GPUTracking/Global/GPUErrorCodes.h +++ b/GPU/GPUTracking/Global/GPUErrorCodes.h @@ -30,3 +30,4 @@ GPUCA_ERROR_CODE(14, ERROR_SLICEDATA_BIN_OVERFLOW) GPUCA_ERROR_CODE(15, ERROR_SLICEDATA_Z_OVERFLOW) GPUCA_ERROR_CODE(16, ERROR_MERGER_HIT_OVERFLOW) GPUCA_ERROR_CODE(17, ERROR_MERGER_TRACK_OVERFLOW) +GPUCA_ERROR_CODE(18, ERROR_COMPRESSION_ROW_HIT_OVERFLOW) From 2c285f5ec418f5cae9a7ed7c46a73c78fcab836c Mon Sep 17 00:00:00 2001 From: David Rohr Date: Thu, 4 Mar 2021 09:58:20 +0100 Subject: [PATCH 18/20] GPU: Minor optimization, GPU can continue processing while we allocate memory --- GPU/GPUTracking/Global/GPUChainTrackingMerger.cxx | 9 +++++---- 1 file changed, 5 insertions(+), 4 deletions(-) diff --git a/GPU/GPUTracking/Global/GPUChainTrackingMerger.cxx b/GPU/GPUTracking/Global/GPUChainTrackingMerger.cxx index 5bbd34c40338e..7ca3fabaa9d57 100644 --- a/GPU/GPUTracking/Global/GPUChainTrackingMerger.cxx +++ b/GPU/GPUTracking/Global/GPUChainTrackingMerger.cxx @@ -273,13 +273,15 @@ int GPUChainTracking::RunTPCTrackingMerger(bool synchronizeOutput) AllocateRegisteredMemory(Merger.MemoryResOutputO2Scratch()); WriteToConstantMemory(RecoStep::TPCMerging, (char*)&processors()->tpcMerger - (char*)processors(), &MergerShadow, sizeof(MergerShadow), 0); runKernel(GetGridAuto(0, deviceType), krnlRunRangeNone, krnlEventNone); - TransferMemoryResourceLinkToHost(RecoStep::TPCMerging, Merger.MemoryResMemory(), 0); - SynchronizeStream(0); + TransferMemoryResourceLinkToHost(RecoStep::TPCMerging, Merger.MemoryResMemory(), 0, &mEvents->single); + runKernel(GetGridAuto(0, deviceType), krnlRunRangeNone, krnlEventNone); + mRec->ReturnVolatileDeviceMemory(); + SynchronizeEvents(&mEvents->single); + ReleaseEvent(&mEvents->single); AllocateRegisteredMemory(Merger.MemoryResOutputO2(), mSubOutputControls[GPUTrackingOutputs::getIndex(&GPUTrackingOutputs::tpcTracksO2)]); AllocateRegisteredMemory(Merger.MemoryResOutputO2Clus(), mSubOutputControls[GPUTrackingOutputs::getIndex(&GPUTrackingOutputs::tpcTracksO2ClusRefs)]); WriteToConstantMemory(RecoStep::TPCMerging, (char*)&processors()->tpcMerger - (char*)processors(), &MergerShadow, sizeof(MergerShadow), 0); - runKernel(GetGridAuto(0, deviceType), krnlRunRangeNone, krnlEventNone); runKernel(GetGridAuto(0, deviceType), krnlRunRangeNone, krnlEventNone); if (GetProcessingSettings().runMC && mIOPtrs.clustersNative && mIOPtrs.clustersNative->clustersMCTruth) { @@ -292,7 +294,6 @@ int GPUChainTracking::RunTPCTrackingMerger(bool synchronizeOutput) TransferMemoryResourceLinkToHost(RecoStep::TPCMerging, Merger.MemoryResOutputO2Clus(), mRec->NStreams() - 2, nullptr, &mEvents->single); ReleaseEvent(&mEvents->single); } - mRec->ReturnVolatileDeviceMemory(); mRec->PopNonPersistentMemory(RecoStep::TPCMerging, qStr2Tag("TPCMERG2")); } #endif From e3cba7f977dfb82d0eb7b071f3cba520bf37b4d0 Mon Sep 17 00:00:00 2001 From: David Rohr Date: Thu, 4 Mar 2021 10:07:05 +0100 Subject: [PATCH 19/20] GPU: Add helper function SynchronizeEventAndRelease --- GPU/GPUTracking/Global/GPUChain.h | 5 +++++ GPU/GPUTracking/Global/GPUChainTrackingCompression.cxx | 6 ++---- GPU/GPUTracking/Global/GPUChainTrackingMerger.cxx | 6 ++---- GPU/GPUTracking/Global/GPUChainTrackingSliceTracker.cxx | 3 +-- 4 files changed, 10 insertions(+), 10 deletions(-) diff --git a/GPU/GPUTracking/Global/GPUChain.h b/GPU/GPUTracking/Global/GPUChain.h index 9758cf3d9ffaa..c3ede06afa7eb 100644 --- a/GPU/GPUTracking/Global/GPUChain.h +++ b/GPU/GPUTracking/Global/GPUChain.h @@ -82,6 +82,11 @@ class GPUChain inline GPUSettingsProcessing& ProcessingSettings() { return mRec->mProcessingSettings; } inline void SynchronizeStream(int stream) { mRec->SynchronizeStream(stream); } inline void SynchronizeEvents(deviceEvent* evList, int nEvents = 1) { mRec->SynchronizeEvents(evList, nEvents); } + inline void SynchronizeEventAndRelease(deviceEvent* ev) + { + SynchronizeEvents(ev); + ReleaseEvent(ev); + } template inline void CondWaitEvent(T& cond, deviceEvent* ev) { diff --git a/GPU/GPUTracking/Global/GPUChainTrackingCompression.cxx b/GPU/GPUTracking/Global/GPUChainTrackingCompression.cxx index fc72c20c22883..d0ad38736d7dd 100644 --- a/GPU/GPUTracking/Global/GPUChainTrackingCompression.cxx +++ b/GPU/GPUTracking/Global/GPUChainTrackingCompression.cxx @@ -52,8 +52,7 @@ int GPUChainTracking::RunTPCCompression() TransferMemoryResourcesToHost(myStep, &Compressor, 0); #ifdef GPUCA_TPC_GEOMETRY_O2 if (mPipelineFinalizationCtx && GetProcessingSettings().doublePipelineClusterizer) { - SynchronizeEvents(&mEvents->single); - ReleaseEvent(&mEvents->single); + SynchronizeEventAndRelease(&mEvents->single); ((GPUChainTracking*)GetNextChainInQueue())->RunTPCClusterizer_prepare(false); ((GPUChainTracking*)GetNextChainInQueue())->mCFContext->ptrClusterNativeSave = processorsShadow()->ioPtrs.clustersNative; } @@ -189,8 +188,7 @@ int GPUChainTracking::RunTPCCompression() } mIOPtrs.tpcCompressedClusters = Compressor.mOutputFlat; if (ProcessingSettings().tpcCompressionGatherMode == 3) { - SynchronizeEvents(&mEvents->stream[outputStream]); - ReleaseEvent(&mEvents->stream[outputStream]); + SynchronizeEventAndRelease(&mEvents->stream[outputStream]); mRec->ReturnVolatileDeviceMemory(); } diff --git a/GPU/GPUTracking/Global/GPUChainTrackingMerger.cxx b/GPU/GPUTracking/Global/GPUChainTrackingMerger.cxx index 7ca3fabaa9d57..96af6eff1fffa 100644 --- a/GPU/GPUTracking/Global/GPUChainTrackingMerger.cxx +++ b/GPU/GPUTracking/Global/GPUChainTrackingMerger.cxx @@ -33,8 +33,7 @@ void GPUChainTracking::RunTPCTrackingMerger_MergeBorderTracks(char withinSlice, runKernel(GetGridAuto(stream, deviceType), krnlRunRangeNone, {nullptr, stream && i < (unsigned int)mRec->NStreams() ? &mEvents->single : nullptr}, i, withinSlice, mergeMode); } ReleaseEvent(&mEvents->single); - SynchronizeEvents(&mEvents->init); - ReleaseEvent(&mEvents->init); + SynchronizeEventAndRelease(&mEvents->init); for (unsigned int i = 0; i < n; i++) { int stream = i % mRec->NStreams(); int n1, n2; @@ -276,8 +275,7 @@ int GPUChainTracking::RunTPCTrackingMerger(bool synchronizeOutput) TransferMemoryResourceLinkToHost(RecoStep::TPCMerging, Merger.MemoryResMemory(), 0, &mEvents->single); runKernel(GetGridAuto(0, deviceType), krnlRunRangeNone, krnlEventNone); mRec->ReturnVolatileDeviceMemory(); - SynchronizeEvents(&mEvents->single); - ReleaseEvent(&mEvents->single); + SynchronizeEventAndRelease(&mEvents->single); AllocateRegisteredMemory(Merger.MemoryResOutputO2(), mSubOutputControls[GPUTrackingOutputs::getIndex(&GPUTrackingOutputs::tpcTracksO2)]); AllocateRegisteredMemory(Merger.MemoryResOutputO2Clus(), mSubOutputControls[GPUTrackingOutputs::getIndex(&GPUTrackingOutputs::tpcTracksO2ClusRefs)]); diff --git a/GPU/GPUTracking/Global/GPUChainTrackingSliceTracker.cxx b/GPU/GPUTracking/Global/GPUChainTrackingSliceTracker.cxx index 45c28bb8433f6..460c56e5242b7 100644 --- a/GPU/GPUTracking/Global/GPUChainTrackingSliceTracker.cxx +++ b/GPU/GPUTracking/Global/GPUChainTrackingSliceTracker.cxx @@ -280,8 +280,7 @@ int GPUChainTracking::RunTPCTrackingSlices_internal() for (int i = 0; i < mRec->NStreams(); i++) { ReleaseEvent(&mEvents->stream[i]); } - SynchronizeEvents(&mEvents->single); - ReleaseEvent(&mEvents->single); + SynchronizeEventAndRelease(&mEvents->single); } if (GetProcessingSettings().debugLevel >= 4) { From 9901dc91b9fa14090dceabb7175e6a50f78eaefb Mon Sep 17 00:00:00 2001 From: David Rohr Date: Thu, 4 Mar 2021 11:18:24 +0100 Subject: [PATCH 20/20] GPU: Add option to free TPC Tracks in O2 format from GPU memory after shipping them to GPU for more memory during compression --- GPU/GPUTracking/Base/GPUReconstruction.cxx | 3 ++ GPU/GPUTracking/Definitions/GPUSettingsList.h | 3 +- .../Global/GPUChainTrackingMerger.cxx | 29 ++++++++++++------- 3 files changed, 23 insertions(+), 12 deletions(-) diff --git a/GPU/GPUTracking/Base/GPUReconstruction.cxx b/GPU/GPUTracking/Base/GPUReconstruction.cxx index 74f42ce24568d..2899af3009373 100644 --- a/GPU/GPUTracking/Base/GPUReconstruction.cxx +++ b/GPU/GPUTracking/Base/GPUReconstruction.cxx @@ -251,6 +251,9 @@ int GPUReconstruction::InitPhaseBeforeDevice() if (mProcessingSettings.createO2Output > 1 && mProcessingSettings.runQA) { mProcessingSettings.createO2Output = 1; } + if (!mProcessingSettings.createO2Output || !IsGPU()) { + mProcessingSettings.clearO2OutputFromGPU = false; + } if (!(mRecoStepsGPU & GPUDataTypes::RecoStep::TPCMerging)) { mProcessingSettings.mergerSortTracks = false; } diff --git a/GPU/GPUTracking/Definitions/GPUSettingsList.h b/GPU/GPUTracking/Definitions/GPUSettingsList.h index e6346aff5d70c..5895791215ada 100644 --- a/GPU/GPUTracking/Definitions/GPUSettingsList.h +++ b/GPU/GPUTracking/Definitions/GPUSettingsList.h @@ -143,7 +143,8 @@ AddOption(rtcConstexpr, bool, true, "", 0, "Replace constant variables by static AddOption(runMC, bool, false, "", 0, "Process MC labels") AddOption(runQA, int, 0, "qa", 'q', "Enable tracking QA (negative number to provide bitmask for QA tasks)", message("Running QA: %s"), def(1)) AddOption(outputSharedClusterMap, bool, false, "", 0, "Ship optional shared cluster map as output for further use") -AddOption(createO2Output, char, 2, "", 0, "Create Track output in O2 format (2 = skip output in GPU format if QA not requested)") +AddOption(createO2Output, char, 2, "", 0, "Create Track output in O2 format (2 = skip non-O2 output in GPU track format (reverts to =1 if QA is requested))") +AddOption(clearO2OutputFromGPU, bool, false, "", 0, "Free the GPU memory used for O2 output after copying to host, prevents further O2 processing on the GPU") AddVariable(eventDisplay, GPUCA_NAMESPACE::gpu::GPUDisplayBackend*, nullptr) AddHelp("help", 'h') EndConfig() diff --git a/GPU/GPUTracking/Global/GPUChainTrackingMerger.cxx b/GPU/GPUTracking/Global/GPUChainTrackingMerger.cxx index 96af6eff1fffa..9cffa0aa17be1 100644 --- a/GPU/GPUTracking/Global/GPUChainTrackingMerger.cxx +++ b/GPU/GPUTracking/Global/GPUChainTrackingMerger.cxx @@ -103,6 +103,7 @@ int GPUChainTracking::RunTPCTrackingMerger(bool synchronizeOutput) GPUTPCGMMerger& Merger = processors()->tpcMerger; GPUTPCGMMerger& MergerShadow = doGPU ? processorsShadow()->tpcMerger : Merger; GPUTPCGMMerger& MergerShadowAll = doGPUall ? processorsShadow()->tpcMerger : Merger; + const int outputStream = mRec->NStreams() - 2; if (GetProcessingSettings().debugLevel >= 2) { GPUInfo("Running TPC Merger"); } @@ -207,7 +208,7 @@ int GPUChainTracking::RunTPCTrackingMerger(bool synchronizeOutput) if (GetProcessingSettings().delayedOutput) { for (unsigned int i = 0; i < mOutputQueue.size(); i++) { - GPUMemCpy(mOutputQueue[i].step, mOutputQueue[i].dst, mOutputQueue[i].src, mOutputQueue[i].size, mRec->NStreams() - 2, false); + GPUMemCpy(mOutputQueue[i].step, mOutputQueue[i].dst, mOutputQueue[i].src, mOutputQueue[i].size, outputStream, false); } mOutputQueue.clear(); } @@ -238,17 +239,17 @@ int GPUChainTracking::RunTPCTrackingMerger(bool synchronizeOutput) if (doGPUall) { RecordMarker(&mEvents->single, 0); if (!GetProcessingSettings().fullMergerOnGPU) { - TransferMemoryResourceLinkToHost(RecoStep::TPCMerging, Merger.MemoryResOutput(), mRec->NStreams() - 2, nullptr, &mEvents->single); + TransferMemoryResourceLinkToHost(RecoStep::TPCMerging, Merger.MemoryResOutput(), outputStream, nullptr, &mEvents->single); } else if (GetProcessingSettings().keepDisplayMemory || GetProcessingSettings().createO2Output <= 1) { - GPUMemCpy(RecoStep::TPCMerging, Merger.OutputTracks(), MergerShadowAll.OutputTracks(), Merger.NOutputTracks() * sizeof(*Merger.OutputTracks()), mRec->NStreams() - 2, 0, nullptr, &mEvents->single); + GPUMemCpy(RecoStep::TPCMerging, Merger.OutputTracks(), MergerShadowAll.OutputTracks(), Merger.NOutputTracks() * sizeof(*Merger.OutputTracks()), outputStream, 0, nullptr, &mEvents->single); if (param().par.dodEdx) { - GPUMemCpy(RecoStep::TPCMerging, Merger.OutputTracksdEdx(), MergerShadowAll.OutputTracksdEdx(), Merger.NOutputTracks() * sizeof(*Merger.OutputTracksdEdx()), mRec->NStreams() - 2, 0, nullptr, &mEvents->single); + GPUMemCpy(RecoStep::TPCMerging, Merger.OutputTracksdEdx(), MergerShadowAll.OutputTracksdEdx(), Merger.NOutputTracks() * sizeof(*Merger.OutputTracksdEdx()), outputStream, 0, nullptr, &mEvents->single); } - GPUMemCpy(RecoStep::TPCMerging, Merger.Clusters(), MergerShadowAll.Clusters(), Merger.NOutputTrackClusters() * sizeof(*Merger.Clusters()), mRec->NStreams() - 2, 0); + GPUMemCpy(RecoStep::TPCMerging, Merger.Clusters(), MergerShadowAll.Clusters(), Merger.NOutputTrackClusters() * sizeof(*Merger.Clusters()), outputStream, 0); if (param().par.earlyTpcTransform) { - GPUMemCpy(RecoStep::TPCMerging, Merger.ClustersXYZ(), MergerShadowAll.ClustersXYZ(), Merger.NOutputTrackClusters() * sizeof(*Merger.ClustersXYZ()), mRec->NStreams() - 2, 0); + GPUMemCpy(RecoStep::TPCMerging, Merger.ClustersXYZ(), MergerShadowAll.ClustersXYZ(), Merger.NOutputTrackClusters() * sizeof(*Merger.ClustersXYZ()), outputStream, 0); } - GPUMemCpy(RecoStep::TPCMerging, Merger.ClusterAttachment(), MergerShadowAll.ClusterAttachment(), Merger.NMaxClusters() * sizeof(*Merger.ClusterAttachment()), mRec->NStreams() - 2, 0); + GPUMemCpy(RecoStep::TPCMerging, Merger.ClusterAttachment(), MergerShadowAll.ClusterAttachment(), Merger.NMaxClusters() * sizeof(*Merger.ClusterAttachment()), outputStream, 0); } ReleaseEvent(&mEvents->single); } else { @@ -277,6 +278,9 @@ int GPUChainTracking::RunTPCTrackingMerger(bool synchronizeOutput) mRec->ReturnVolatileDeviceMemory(); SynchronizeEventAndRelease(&mEvents->single); + if (GetProcessingSettings().clearO2OutputFromGPU) { + mRec->AllocateVolatileDeviceMemory(0); // make future device memory allocation volatile + } AllocateRegisteredMemory(Merger.MemoryResOutputO2(), mSubOutputControls[GPUTrackingOutputs::getIndex(&GPUTrackingOutputs::tpcTracksO2)]); AllocateRegisteredMemory(Merger.MemoryResOutputO2Clus(), mSubOutputControls[GPUTrackingOutputs::getIndex(&GPUTrackingOutputs::tpcTracksO2ClusRefs)]); WriteToConstantMemory(RecoStep::TPCMerging, (char*)&processors()->tpcMerger - (char*)processors(), &MergerShadow, sizeof(MergerShadow), 0); @@ -288,15 +292,18 @@ int GPUChainTracking::RunTPCTrackingMerger(bool synchronizeOutput) runKernel(GetGridAuto(0, GPUReconstruction::krnlDeviceType::CPU), krnlRunRangeNone, krnlEventNone); } else { RecordMarker(&mEvents->single, 0); - TransferMemoryResourceLinkToHost(RecoStep::TPCMerging, Merger.MemoryResOutputO2(), mRec->NStreams() - 2, nullptr, &mEvents->single); - TransferMemoryResourceLinkToHost(RecoStep::TPCMerging, Merger.MemoryResOutputO2Clus(), mRec->NStreams() - 2, nullptr, &mEvents->single); + TransferMemoryResourceLinkToHost(RecoStep::TPCMerging, Merger.MemoryResOutputO2(), outputStream, nullptr, &mEvents->single); + TransferMemoryResourceLinkToHost(RecoStep::TPCMerging, Merger.MemoryResOutputO2Clus(), outputStream, nullptr, &mEvents->single); ReleaseEvent(&mEvents->single); } mRec->PopNonPersistentMemory(RecoStep::TPCMerging, qStr2Tag("TPCMERG2")); } #endif - if (synchronizeOutput) { - SynchronizeStream(mRec->NStreams() - 2); + if (synchronizeOutput || GetProcessingSettings().clearO2OutputFromGPU) { + SynchronizeStream(outputStream); + } + if (GetProcessingSettings().clearO2OutputFromGPU) { + mRec->ReturnVolatileDeviceMemory(); } mIOPtrs.mergedTracks = Merger.OutputTracks();