diff --git a/Detectors/Base/include/DetectorsBase/Propagator.h b/Detectors/Base/include/DetectorsBase/Propagator.h index c9b71b0d123a6..72fa4032847b6 100644 --- a/Detectors/Base/include/DetectorsBase/Propagator.h +++ b/Detectors/Base/include/DetectorsBase/Propagator.h @@ -111,6 +111,7 @@ class Propagator GPUd() const o2::base::MatLayerCylSet* getMatLUT() const { return mMatLUT; } GPUd() void setGPUField(const o2::gpu::GPUTPCGMPolynomialField* field) { mGPUField = field; } GPUd() const o2::gpu::GPUTPCGMPolynomialField* getGPUField() const { return mGPUField; } + GPUd() void setBz(float bz) { mBz = bz; } #ifndef GPUCA_GPUCODE static Propagator* Instance() diff --git a/Detectors/Base/src/Propagator.cxx b/Detectors/Base/src/Propagator.cxx index 951f014052b80..098d3bc440376 100644 --- a/Detectors/Base/src/Propagator.cxx +++ b/Detectors/Base/src/Propagator.cxx @@ -98,7 +98,10 @@ int Propagator::initFieldFromGRP(const o2::parameters::GRPObject* grp, bool verb } return 0; } - +#elif !defined(GPUCA_GPUCODE) +Propagator::Propagator() +{ +} // empty dummy constructor for standalone benchmark #endif //_______________________________________________________________________ diff --git a/Detectors/TPC/workflow/src/CATrackerSpec.cxx b/Detectors/TPC/workflow/src/CATrackerSpec.cxx index 297b436e4c2d3..86833d421a0a8 100644 --- a/Detectors/TPC/workflow/src/CATrackerSpec.cxx +++ b/Detectors/TPC/workflow/src/CATrackerSpec.cxx @@ -39,6 +39,9 @@ #include "TPCdEdxCalibrationSplines.h" #include "DPLUtils/DPLRawParser.h" #include "DetectorsBase/MatLayerCylSet.h" +#include "DetectorsBase/Propagator.h" +#include "DetectorsBase/GeometryManager.h" +#include "DetectorsCommonDataFormats/NameConf.h" #include "DetectorsRaw/HBFUtils.h" #include "TPCBase/RDHUtils.h" #include "GPUO2InterfaceConfiguration.h" @@ -127,7 +130,9 @@ DataProcessorSpec getCATrackerSpec(CompletionPolicyData* policyData, ca::Config tracker = std::make_unique(); // Create configuration object and fill settings - const auto grp = o2::parameters::GRPObject::loadFrom("o2sim_grp.root"); + const auto grp = o2::parameters::GRPObject::loadFrom(o2::base::NameConf::getGRPFileName()); + o2::base::GeometryManager::loadGeometry(); + o2::base::Propagator::initFieldFromGRP(o2::base::NameConf::getGRPFileName()); if (grp) { config.configEvent.solenoidBz = 5.00668f * grp->getL3Current() / 30000.; config.configEvent.continuousMaxTimeBin = grp->isDetContinuousReadOut(o2::detectors::DetID::TPC) ? -1 : 0; // Number of timebins in timeframe if continuous, 0 otherwise @@ -216,15 +221,16 @@ DataProcessorSpec getCATrackerSpec(CompletionPolicyData* policyData, ca::Config if (config.configCalib.fastTransform == nullptr) { throw std::invalid_argument("GPUCATracking: initialization of the TPC transformation failed"); } + if (confParam.matLUTFile.size()) { config.configCalib.matLUT = o2::base::MatLayerCylSet::loadFromFile(confParam.matLUTFile.c_str(), "MatBud"); } + if (confParam.dEdxFile.size()) { processAttributes->dEdxSplines.reset(new TPCdEdxCalibrationSplines(confParam.dEdxFile.c_str())); } else { processAttributes->dEdxSplines.reset(new TPCdEdxCalibrationSplines); } - config.configCalib.dEdxSplines = processAttributes->dEdxSplines.get(); if (boost::filesystem::exists(confParam.gainCalibFile)) { @@ -239,9 +245,10 @@ DataProcessorSpec getCATrackerSpec(CompletionPolicyData* policyData, ca::Config } config.configCalib.tpcPadGain = processAttributes->tpcPadGainCalib.get(); + config.configCalib.o2Propagator = Propagator::Instance(); + // Sample code what needs to be done for the TRD Geometry, when we extend this to TRD tracking. - /*o2::base::GeometryManager::loadGeometry(); - o2::trd::Geometry gm; + /* o2::trd::Geometry gm; gm.createPadPlaneArray(); gm.createClusterMatrixArray(); std::unique_ptr gf(gm); diff --git a/GPU/Common/GPUDefGPUParameters.h b/GPU/Common/GPUDefGPUParameters.h index 43c19f42fa2e2..19ca5e8df42ba 100644 --- a/GPU/Common/GPUDefGPUParameters.h +++ b/GPU/Common/GPUDefGPUParameters.h @@ -345,6 +345,12 @@ #ifndef GPUCA_LB_GPUTPCStartHitsSorter #define GPUCA_LB_GPUTPCStartHitsSorter 256 #endif + #ifndef GPUCA_LB_GPUTrackingRefitKernel_mode0asGPU + #define GPUCA_LB_GPUTrackingRefitKernel_mode0asGPU 256 + #endif + #ifndef GPUCA_LB_GPUTrackingRefitKernel_mode1asTrackParCov + #define GPUCA_LB_GPUTrackingRefitKernel_mode1asTrackParCov 256 + #endif #define GPUCA_GET_THREAD_COUNT(...) GPUCA_M_FIRST(__VA_ARGS__) #else // The following defaults are needed to compile the host code diff --git a/GPU/GPUTracking/Base/GPUConstantMem.h b/GPU/GPUTracking/Base/GPUConstantMem.h index d498be62793b1..8372eaa7a9582 100644 --- a/GPU/GPUTracking/Base/GPUConstantMem.h +++ b/GPU/GPUTracking/Base/GPUConstantMem.h @@ -47,6 +47,7 @@ class GPUTRDTracker_t #include "GPUTPCCompression.h" #include "GPUITSFitter.h" #include "GPUTPCClusterFinder.h" +#include "GPUTrackingRefit.h" #else #include "GPUO2FakeClasses.h" #endif @@ -71,6 +72,7 @@ struct GPUConstantMem { GPUTRDTrackerGPU trdTracker; GPUTPCClusterFinder tpcClusterer[GPUCA_NSLICES]; GPUITSFitter itsFitter; + GPUTrackingRefitProcessor trackingRefit; GPUTrackingInOutPointers ioPtrs; GPUCalibObjectsConst calibObjects; GPUErrors errorCodes; @@ -133,13 +135,6 @@ GPUdi() void GPUProcessor::raiseError(unsigned int code, unsigned int param1, un GetConstantMem()->errorCodes.raiseError(code, param1, param2, param3); } -#if defined(GPUCA_NOCOMPAT_ALLCINT) && (!defined(GPUCA_GPULIBRARY) || !defined(GPUCA_ALIROOT_LIB)) && defined(HAVE_O2HEADERS) -GPUd() float GPUTPCClusterFinder::getGainCorrection(tpccf::Row row, tpccf::Pad pad) const -{ - return GetConstantMem()->calibObjects.tpcPadGain->getGainCorrection(mISlice, row, pad); -} -#endif - } // namespace gpu } // namespace GPUCA_NAMESPACE diff --git a/GPU/GPUTracking/Base/GPUDataTypes.h b/GPU/GPUTracking/Base/GPUDataTypes.h index c948f43e8bbe0..6612f77ee229c 100644 --- a/GPU/GPUTracking/Base/GPUDataTypes.h +++ b/GPU/GPUTracking/Base/GPUDataTypes.h @@ -48,6 +48,7 @@ namespace o2 class MCCompLabel; namespace base { +class Propagator; class MatLayerCylSet; } // namespace base namespace trd @@ -136,7 +137,7 @@ class GPUDataTypes TPCRaw = 64 }; #ifdef GPUCA_NOCOMPAT_ALLOPENCL - static constexpr const char* const RECO_STEP_NAMES[] = {"TPC Transformation", "TPC Sector Tracking", "TPC Track Merging and Fit", "TPC Compression", "TRD Tracking", "ITS Tracking", "TPC dEdx Computation", "TPC Cluster Finding", "TPC Decompression"}; + static constexpr const char* const RECO_STEP_NAMES[] = {"TPC Transformation", "TPC Sector Tracking", "TPC Track Merging and Fit", "TPC Compression", "TRD Tracking", "ITS Tracking", "TPC dEdx Computation", "TPC Cluster Finding", "TPC Decompression", "Global Refit"}; static constexpr const char* const GENERAL_STEP_NAMES[] = {"Prepare", "QA"}; typedef bitfield RecoStepField; typedef bitfield InOutTypeField; @@ -176,6 +177,7 @@ struct GPUCalibObjectsTemplate { typename S::type* trdGeometry = nullptr; typename S::type* dEdxSplines = nullptr; typename S::type* tpcPadGain = nullptr; + typename S::type* o2Propagator = nullptr; }; typedef GPUCalibObjectsTemplate GPUCalibObjects; typedef GPUCalibObjectsTemplate GPUCalibObjectsConst; diff --git a/GPU/GPUTracking/Base/GPUO2DataTypes.h b/GPU/GPUTracking/Base/GPUO2DataTypes.h index 40d814e6a764d..66a53883cb8fc 100644 --- a/GPU/GPUTracking/Base/GPUO2DataTypes.h +++ b/GPU/GPUTracking/Base/GPUO2DataTypes.h @@ -20,6 +20,7 @@ #include "DataFormatsTPC/ClusterNative.h" #include "DataFormatsTPC/Digit.h" #include "DetectorsBase/MatLayerCylSet.h" +#include "DetectorsBase/Propagator.h" #include "TRDBase/GeometryFlat.h" #else #include "GPUO2FakeClasses.h" diff --git a/GPU/GPUTracking/Base/GPUO2FakeClasses.h b/GPU/GPUTracking/Base/GPUO2FakeClasses.h index 68c592366fd76..5449549c2a25e 100644 --- a/GPU/GPUTracking/Base/GPUO2FakeClasses.h +++ b/GPU/GPUTracking/Base/GPUO2FakeClasses.h @@ -102,6 +102,9 @@ class GPUTPCCompression class GPUTPCClusterFinder { }; +class GPUTrackingRefitProcessor +{ +}; struct GPUTPCCFChainContext { }; #ifndef __OPENCL__ diff --git a/GPU/GPUTracking/Base/GPUParam.cxx b/GPU/GPUTracking/Base/GPUParam.cxx index a338bc8202a6b..c8292f65fcfe8 100644 --- a/GPU/GPUTracking/Base/GPUParam.cxx +++ b/GPU/GPUTracking/Base/GPUParam.cxx @@ -27,6 +27,9 @@ using namespace GPUCA_NAMESPACE::gpu; #endif #include #include +#ifdef HAVE_O2HEADERS +#include "DetectorsBase/Propagator.h" +#endif #include "utils/qconfigrtc.h" @@ -244,3 +247,21 @@ std::string GPUParamRTC::generateRTCCode(const GPUParam& param, bool useConstexp static_assert(alignof(GPUCA_NAMESPACE::gpu::GPUParam) == alignof(GPUCA_NAMESPACE::gpu::GPUSettingsRec)); static_assert(alignof(GPUCA_NAMESPACE::gpu::GPUParam) == alignof(GPUCA_NAMESPACE::gpu::GPUSettingsParam)); static_assert(sizeof(GPUCA_NAMESPACE::gpu::GPUParam) - sizeof(GPUCA_NAMESPACE::gpu::GPUParamRTC) == sizeof(GPUCA_NAMESPACE::gpu::GPUSettingsRec) + sizeof(GPUCA_NAMESPACE::gpu::GPUSettingsParam) - sizeof(GPUCA_NAMESPACE::gpu::gpu_rtc::GPUSettingsRec) - sizeof(GPUCA_NAMESPACE::gpu::gpu_rtc::GPUSettingsParam)); + +o2::base::Propagator* GPUParam::GetDefaultO2Propagator(bool useGPUField) const +{ + o2::base::Propagator* prop = nullptr; +#ifdef HAVE_O2HEADERS + if (useGPUField == false) { + throw std::runtime_error("o2 propagator withouzt gpu field unsupported"); + } + prop = o2::base::Propagator::Instance(); + if (useGPUField) { + prop->setGPUField(&polynomialField); + prop->setBz(polynomialField.GetNominalBz()); + } +#else + throw std::runtime_error("o2 propagator unsupported"); +#endif + return prop; +} diff --git a/GPU/GPUTracking/Base/GPUParam.h b/GPU/GPUTracking/Base/GPUParam.h index c1dc43d392e0b..031cc8b78eee6 100644 --- a/GPU/GPUTracking/Base/GPUParam.h +++ b/GPU/GPUTracking/Base/GPUParam.h @@ -21,6 +21,14 @@ #include "GPUTPCGeometry.h" #include "GPUTPCGMPolynomialField.h" +namespace o2 +{ +namespace base +{ +class Propagator; +} // namespace base +} // namespace o2 + namespace GPUCA_NAMESPACE { namespace gpu @@ -63,6 +71,7 @@ struct GPUParam : public internal::GPUParam_t void SetDefaults(const GPUSettingsEvent* e, const GPUSettingsRec* r = nullptr, const GPUSettingsProcessing* p = nullptr, const GPURecoStepConfiguration* w = nullptr); void UpdateEventSettings(const GPUSettingsEvent* e, const GPUSettingsProcessing* p = nullptr); void LoadClusterErrors(bool Print = 0); + o2::base::Propagator* GetDefaultO2Propagator(bool useGPUField = false) const; #endif GPUd() float Alpha(int iSlice) const diff --git a/GPU/GPUTracking/Base/GPUReconstructionCPU.h b/GPU/GPUTracking/Base/GPUReconstructionCPU.h index 8b0bc183d70a1..3b9866f361f0e 100644 --- a/GPU/GPUTracking/Base/GPUReconstructionCPU.h +++ b/GPU/GPUTracking/Base/GPUReconstructionCPU.h @@ -39,6 +39,7 @@ #include "GPUTPCConvertKernel.h" #include "GPUTPCCompressionKernels.h" #include "GPUTPCClusterFinderKernels.h" +#include "GPUTrackingRefitKernel.h" #endif namespace GPUCA_NAMESPACE diff --git a/GPU/GPUTracking/Base/GPUReconstructionIncludesDevice.h b/GPU/GPUTracking/Base/GPUReconstructionIncludesDevice.h index 6ae8c9119c467..f05f120c1b18b 100644 --- a/GPU/GPUTracking/Base/GPUReconstructionIncludesDevice.h +++ b/GPU/GPUTracking/Base/GPUReconstructionIncludesDevice.h @@ -97,6 +97,10 @@ using namespace GPUCA_NAMESPACE::gpu; // Files for ITS Track Fit #include "GPUITSFitterKernels.cxx" +// Files for Refit +#include "GPUTrackingRefit.cxx" +#include "GPUTrackingRefitKernel.cxx" + #if !defined(GPUCA_O2_LIB) && defined(__HIPCC__) && !defined(GPUCA_NO_ITS_TRAITS) && !defined(GPUCA_GPUCODE_GENRTC) #include "VertexerTraitsHIP.hip.cxx" #include "ContextHIP.hip.cxx" diff --git a/GPU/GPUTracking/Base/GPUReconstructionKernels.h b/GPU/GPUTracking/Base/GPUReconstructionKernels.h index 904b798debb6d..c300714ccf3a0 100644 --- a/GPU/GPUTracking/Base/GPUReconstructionKernels.h +++ b/GPU/GPUTracking/Base/GPUReconstructionKernels.h @@ -69,25 +69,26 @@ GPUCA_KRNL_LB((GPUTPCCompressionGatherKernels, buffered32 ), (simple), (), ()) GPUCA_KRNL_LB((GPUTPCCompressionGatherKernels, buffered64 ), (simple), (), ()) GPUCA_KRNL_LB((GPUTPCCompressionGatherKernels, buffered128 ), (simple), (), ()) GPUCA_KRNL_LB((GPUTPCCompressionGatherKernels, multiBlock ), (simple), (), ()) - GPUCA_KRNL_LB((GPUTPCCFCheckPadBaseline ), (single), (), ()) -GPUCA_KRNL_LB((GPUTPCCFChargeMapFiller, fillIndexMap ), (single), (), ()) -GPUCA_KRNL_LB((GPUTPCCFChargeMapFiller, fillFromDigits ), (single), (), ()) -GPUCA_KRNL_LB((GPUTPCCFChargeMapFiller, findFragmentStart ), (single), (, char setPositions), (, setPositions)) +GPUCA_KRNL_LB((GPUTPCCFChargeMapFiller, fillIndexMap ), (single), (), ()) +GPUCA_KRNL_LB((GPUTPCCFChargeMapFiller, fillFromDigits ), (single), (), ()) +GPUCA_KRNL_LB((GPUTPCCFChargeMapFiller, findFragmentStart ), (single), (, char setPositions), (, setPositions)) GPUCA_KRNL_LB((GPUTPCCFPeakFinder ), (single), (), ()) -GPUCA_KRNL_LB((GPUTPCCFNoiseSuppression, noiseSuppression ), (single), (), ()) -GPUCA_KRNL_LB((GPUTPCCFNoiseSuppression, updatePeaks ), (single), (), ()) +GPUCA_KRNL_LB((GPUTPCCFNoiseSuppression, noiseSuppression ), (single), (), ()) +GPUCA_KRNL_LB((GPUTPCCFNoiseSuppression, updatePeaks ), (single), (), ()) GPUCA_KRNL_LB((GPUTPCCFDeconvolution ), (single), (), ()) GPUCA_KRNL_LB((GPUTPCCFClusterizer ), (single), (, char onlyMC), (, onlyMC)) -GPUCA_KRNL(( GPUTPCCFMCLabelFlattener, setRowOffsets ), (single), (), ()) -GPUCA_KRNL(( GPUTPCCFMCLabelFlattener, flatten ), (single), (, GPUPtr1(GPUTPCLinearLabels*, out)), (, GPUPtr2(GPUTPCLinearLabels*, out))) -GPUCA_KRNL_LB((GPUTPCCFStreamCompaction, scanStart ), (single), (, int iBuf, int stage), (, iBuf, stage)) -GPUCA_KRNL_LB((GPUTPCCFStreamCompaction, scanUp ), (single), (, int iBuf, int nElems), (, iBuf, nElems)) -GPUCA_KRNL_LB((GPUTPCCFStreamCompaction, scanTop ), (single), (, int iBuf, int nElems), (, iBuf, nElems)) -GPUCA_KRNL_LB((GPUTPCCFStreamCompaction, scanDown ), (single), (, int iBuf, unsigned int offset, int nElems), (, iBuf, offset, nElems)) -GPUCA_KRNL_LB((GPUTPCCFStreamCompaction, compactDigits ), (single), (, int iBuf, int stage, GPUPtr1(ChargePos*, in), GPUPtr1(ChargePos*, out)), (, iBuf, stage, GPUPtr2(ChargePos*, in), GPUPtr2(ChargePos*, out))) +GPUCA_KRNL(( GPUTPCCFMCLabelFlattener, setRowOffsets ), (single), (), ()) +GPUCA_KRNL(( GPUTPCCFMCLabelFlattener, flatten ), (single), (, GPUPtr1(GPUTPCLinearLabels*, out)), (, GPUPtr2(GPUTPCLinearLabels*, out))) +GPUCA_KRNL_LB((GPUTPCCFStreamCompaction, scanStart ), (single), (, int iBuf, int stage), (, iBuf, stage)) +GPUCA_KRNL_LB((GPUTPCCFStreamCompaction, scanUp ), (single), (, int iBuf, int nElems), (, iBuf, nElems)) +GPUCA_KRNL_LB((GPUTPCCFStreamCompaction, scanTop ), (single), (, int iBuf, int nElems), (, iBuf, nElems)) +GPUCA_KRNL_LB((GPUTPCCFStreamCompaction, scanDown ), (single), (, int iBuf, unsigned int offset, int nElems), (, iBuf, offset, nElems)) +GPUCA_KRNL_LB((GPUTPCCFStreamCompaction, compactDigits ), (single), (, int iBuf, int stage, GPUPtr1(ChargePos*, in), GPUPtr1(ChargePos*, out)), (, iBuf, stage, GPUPtr2(ChargePos*, in), GPUPtr2(ChargePos*, out))) GPUCA_KRNL_LB((GPUTPCCFDecodeZS ), (single), (, int firstHBF), (, firstHBF)) GPUCA_KRNL_LB((GPUTPCCFGather ), (single), (, GPUPtr1(o2::tpc::ClusterNative*, dest)), (, GPUPtr2(o2::tpc::ClusterNative*, dest))) +GPUCA_KRNL_LB((GPUTrackingRefitKernel, mode0asGPU ), (simple), (), ()) +GPUCA_KRNL_LB((GPUTrackingRefitKernel, mode1asTrackParCov ), (simple), (), ()) #endif #endif // clang-format on diff --git a/GPU/GPUTracking/Base/GPUSettingsList.h b/GPU/GPUTracking/Base/GPUSettingsList.h index 6b385d6d97c56..146996d50a245 100644 --- a/GPU/GPUTracking/Base/GPUSettingsList.h +++ b/GPU/GPUTracking/Base/GPUSettingsList.h @@ -85,6 +85,7 @@ AddOptionRTC(loopInterpolationInExtraPass, char, -1, "", 0, "Perform loop interp AddOptionRTC(mergerReadFromTrackerDirectly, char, 1, "", 0, "Forward data directly from tracker to merger on GPU") AddOptionRTC(useMatLUT, char, 0, "", 0, "Use material lookup table for TPC refit") AddOptionRTC(trdStopTrkAfterNMissLy, unsigned char, 6, "", 0, "Abandon track following after N layers without a TRD match") +AddOptionRTC(trackingRefitGPUModel, char, 1, "", 0, "Use GPU track model for the Global Track Refit") AddCustomCPP(void SetMinTrackPt(float v) { MaxTrackQPt = v > 0.001 ? (1. / v) : (1. / 0.001); }) AddVariable(dummyRTC, float, 0.f) // Ensure non empty struct and proper alignment even if all normal members are constexpr AddHelp("help", 'h') diff --git a/GPU/GPUTracking/CMakeLists.txt b/GPU/GPUTracking/CMakeLists.txt index 434cff83d29ba..7f340f56467e2 100644 --- a/GPU/GPUTracking/CMakeLists.txt +++ b/GPU/GPUTracking/CMakeLists.txt @@ -175,7 +175,8 @@ if(ALIGPU_BUILD_TYPE STREQUAL "O2" OR CONFIG_O2_EXTENSIONS) TPCClusterFinder/GPUTPCCFGather.cxx TPCClusterFinder/TPCPadGainCalib.cxx dEdx/TPCdEdxCalibrationSplines.cxx - Refit/GPUTrackingRefit.cxx) + Refit/GPUTrackingRefit.cxx + Refit/GPUTrackingRefitKernel.cxx) set(SRCS_NO_H ${SRCS_NO_H} TPCClusterFinder/GPUTPCClusterFinderDump.cxx) diff --git a/GPU/GPUTracking/Global/GPUChainTracking.cxx b/GPU/GPUTracking/Global/GPUChainTracking.cxx index 55e26c31d8133..863bc98e878b2 100644 --- a/GPU/GPUTracking/Global/GPUChainTracking.cxx +++ b/GPU/GPUTracking/Global/GPUChainTracking.cxx @@ -125,6 +125,9 @@ void GPUChainTracking::RegisterPermanentMemoryAndProcessors() mRec->RegisterGPUProcessor(&processors()->tpcClusterer[i], GetRecoStepsGPU() & RecoStep::TPCClusterFinding); } } + if (GetRecoSteps() & RecoStep::Refit) { + mRec->RegisterGPUProcessor(&processors()->trackingRefit, GetRecoStepsGPU() & RecoStep::Refit); + } #endif #ifdef GPUCA_KERNEL_DEBUGGER_OUTPUT mRec->RegisterGPUProcessor(&processors()->debugOutput, true); @@ -162,6 +165,9 @@ void GPUChainTracking::RegisterGPUProcessors() mRec->RegisterGPUDeviceProcessor(&processorsShadow()->tpcClusterer[i], &processors()->tpcClusterer[i]); } } + if (GetRecoStepsGPU() & RecoStep::Refit) { + mRec->RegisterGPUDeviceProcessor(&processorsShadow()->trackingRefit, &processors()->trackingRefit); + } #endif #ifdef GPUCA_KERNEL_DEBUGGER_OUTPUT mRec->RegisterGPUDeviceProcessor(&processorsShadow()->debugOutput, &processors()->debugOutput); @@ -259,6 +265,10 @@ bool GPUChainTracking::ValidateSteps() GPUError("Cannot run gain calibration without calibration object"); return false; } + if ((GetRecoSteps() & GPUDataTypes::RecoStep::Refit) && !param().rec.trackingRefitGPUModel && (processors()->calibObjects.o2Propagator == nullptr || processors()->calibObjects.matLUT == nullptr)) { + GPUError("Cannot run refit with o2 track model without o2 propagator"); + return false; + } return true; } @@ -386,8 +396,15 @@ int GPUChainTracking::Init() if (processors()->calibObjects.tpcPadGain) { memcpy((void*)mFlatObjectsShadow.mCalibObjects.tpcPadGain, (const void*)processors()->calibObjects.tpcPadGain, sizeof(*processors()->calibObjects.tpcPadGain)); } + if (processors()->calibObjects.o2Propagator) { + memcpy((void*)mFlatObjectsShadow.mCalibObjects.o2Propagator, (const void*)processors()->calibObjects.o2Propagator, sizeof(*processors()->calibObjects.o2Propagator)); + mFlatObjectsShadow.mCalibObjects.o2Propagator->setGPUField(&processorsDevice()->param.polynomialField); + mFlatObjectsShadow.mCalibObjects.o2Propagator->setBz(param().polynomialField.GetNominalBz()); + mFlatObjectsShadow.mCalibObjects.o2Propagator->setMatLUT(mFlatObjectsShadow.mCalibObjects.matLUT); + } #endif TransferMemoryResourceLinkToGPU(RecoStep::NoRecoStep, mFlatObjectsShadow.mMemoryResFlat); + memcpy((void*)&processorsShadow()->calibObjects, (void*)&mFlatObjectsDevice.mCalibObjects, sizeof(mFlatObjectsDevice.mCalibObjects)); WriteToConstantMemory(RecoStep::NoRecoStep, (char*)&processors()->calibObjects - (char*)processors(), &mFlatObjectsDevice.mCalibObjects, sizeof(mFlatObjectsDevice.mCalibObjects), -1); // First initialization, for users not using RunChain processorsShadow()->errorCodes.setMemory(mInputsShadow->mErrorCodes); WriteToConstantMemory(RecoStep::NoRecoStep, (char*)&processors()->errorCodes - (char*)processors(), &processorsShadow()->errorCodes, sizeof(processorsShadow()->errorCodes), -1); @@ -454,7 +471,9 @@ void* GPUChainTracking::GPUTrackingFlatObjects::SetPointersFlatObjects(void* mem if (mChainTracking->GetTRDGeometry()) { computePointerWithAlignment(mem, mCalibObjects.trdGeometry, 1); } - + if (mChainTracking->GetO2Propagator()) { + computePointerWithAlignment(mem, mCalibObjects.o2Propagator, 1); + } #endif return mem; } @@ -2393,22 +2412,24 @@ int GPUChainTracking::DoTRDGPUTracking() int GPUChainTracking::RunRefit() { #ifdef HAVE_O2HEADERS - GPUTrackingRefit re; - re.SetPtrsFromGPUConstantMem(processorsShadow()); - re.SetPropagatorDefault(); - for (unsigned int i = 0; i < mIOPtrs.nMergedTracks; i++) { - if (mIOPtrs.mergedTracks[i].OK()) { - printf("\nRefitting track %d\n", i); - GPUTPCGMMergedTrack t = mIOPtrs.mergedTracks[i]; - int retval = re.RefitTrackAsGPU(t, false, true); - printf("Refit error code: %d\n", retval); + bool doGPU = GetRecoStepsGPU() & RecoStep::Refit; + GPUTrackingRefitProcessor& Refit = processors()->trackingRefit; + GPUTrackingRefitProcessor& RefitShadow = doGPU ? processorsShadow()->trackingRefit : Refit; - printf("\nRefitting track TrackParCov %d\n", i); - t = mIOPtrs.mergedTracks[i]; - retval = re.RefitTrackAsTrackParCov(t, false, true); - printf("Refit error code: %d\n", retval); - } + const auto& threadContext = GetThreadContext(); + SetupGPUProcessor(&Refit, false); + RefitShadow.SetPtrsFromGPUConstantMem(processorsShadow(), doGPU ? &processorsDevice()->param : nullptr); + RefitShadow.SetPropagator(doGPU ? processorsShadow()->calibObjects.o2Propagator : GetO2Propagator()); + RefitShadow.mPTracks = (doGPU ? processorsShadow() : processors())->tpcMerger.OutputTracks(); + WriteToConstantMemory(RecoStep::Refit, (char*)&processors()->trackingRefit - (char*)processors(), &RefitShadow, sizeof(RefitShadow), 0); + //TransferMemoryResourcesToGPU(RecoStep::Refit, &Refit, 0); + if (param().rec.trackingRefitGPUModel) { + runKernel(GetGrid(mIOPtrs.nMergedTracks, 0), krnlRunRangeNone); + } else { + runKernel(GetGrid(mIOPtrs.nMergedTracks, 0), krnlRunRangeNone); } + //TransferMemoryResourcesToHost(RecoStep::Refit, &Refit, 0); + SynchronizeStream(0); #endif return 0; } @@ -2657,3 +2678,10 @@ void GPUChainTracking::ClearErrorCodes() } TransferMemoryResourceLinkToGPU(RecoStep::NoRecoStep, mInputsHost->mResourceErrorCodes, 0); } + +void GPUChainTracking::SetDefaultO2PropagatorForGPU() +{ + o2::base::Propagator* prop = param().GetDefaultO2Propagator(true); + prop->setMatLUT(processors()->calibObjects.matLUT); + SetO2Propagator(prop); +} diff --git a/GPU/GPUTracking/Global/GPUChainTracking.h b/GPU/GPUTracking/Global/GPUChainTracking.h index 3d93cb85a080d..a757e3598934a 100644 --- a/GPU/GPUTracking/Global/GPUChainTracking.h +++ b/GPU/GPUTracking/Global/GPUChainTracking.h @@ -156,6 +156,7 @@ class GPUChainTracking : public GPUChain, GPUReconstructionHelpers::helperDelega const TPCdEdxCalibrationSplines* GetdEdxSplines() const { return processors()->calibObjects.dEdxSplines; } const o2::base::MatLayerCylSet* GetMatLUT() const { return processors()->calibObjects.matLUT; } const GPUTRDGeometry* GetTRDGeometry() const { return (GPUTRDGeometry*)processors()->calibObjects.trdGeometry; } + const o2::base::Propagator* GetO2Propagator() const { return processors()->calibObjects.o2Propagator; } void SetTPCFastTransform(std::unique_ptr&& tpcFastTransform); void SetdEdxSplines(std::unique_ptr&& dEdxSplines); void SetMatLUT(std::unique_ptr&& lut); @@ -165,6 +166,8 @@ class GPUChainTracking : public GPUChain, GPUReconstructionHelpers::helperDelega void SetdEdxSplines(const TPCdEdxCalibrationSplines* dEdxSplines) { processors()->calibObjects.dEdxSplines = dEdxSplines; } void SetMatLUT(const o2::base::MatLayerCylSet* lut) { processors()->calibObjects.matLUT = lut; } void SetTRDGeometry(const o2::trd::GeometryFlat* geo) { processors()->calibObjects.trdGeometry = geo; } + void SetO2Propagator(const o2::base::Propagator* prop) { processors()->calibObjects.o2Propagator = prop; } + void SetDefaultO2PropagatorForGPU(); void LoadClusterErrors(); void SetOutputControlCompressedClusters(GPUOutputControl* v) { mOutputCompressedClusters = v; } void SetOutputControlClustersNative(GPUOutputControl* v) { mOutputClustersNative = v; } diff --git a/GPU/GPUTracking/Interface/GPUO2Interface.cxx b/GPU/GPUTracking/Interface/GPUO2Interface.cxx index d5df395ed5d21..40ea66aa7157e 100644 --- a/GPU/GPUTracking/Interface/GPUO2Interface.cxx +++ b/GPU/GPUTracking/Interface/GPUO2Interface.cxx @@ -57,6 +57,7 @@ int GPUTPCO2Interface::Initialize(const GPUO2InterfaceConfiguration& config) mChain->SetdEdxSplines(mConfig->configCalib.dEdxSplines); mChain->SetMatLUT(mConfig->configCalib.matLUT); mChain->SetTRDGeometry(mConfig->configCalib.trdGeometry); + mChain->SetO2Propagator(mConfig->configCalib.o2Propagator); if (mConfig->configInterface.outputToExternalBuffers) { mOutputCompressedClusters.reset(new GPUOutputControl); mChain->SetOutputControlCompressedClusters(mOutputCompressedClusters.get()); diff --git a/GPU/GPUTracking/Refit/GPUTrackingRefit.cxx b/GPU/GPUTracking/Refit/GPUTrackingRefit.cxx index 7d86bd11e0f04..809bc6929f16e 100644 --- a/GPU/GPUTracking/Refit/GPUTrackingRefit.cxx +++ b/GPU/GPUTracking/Refit/GPUTrackingRefit.cxx @@ -24,14 +24,14 @@ #include "DetectorsBase/Propagator.h" #include "DataFormatsTPC/TrackTPC.h" #include "GPUParam.inc" +#include "GPUCommonArray.h" +#include "GPUParam.h" using namespace GPUCA_NAMESPACE::gpu; using namespace o2::track; using namespace o2::base; using namespace o2::tpc; -static constexpr float kDeg2Rad = M_PI / 180.f; -static constexpr float kSectAngle = 2 * M_PI / 18.f; static constexpr int kIGNORE_ENDS = 3; #define IgnoreErrors(SNP) \ @@ -45,6 +45,7 @@ static constexpr int kIGNORE_ENDS = 3; } // End IgnoreErrors +#ifndef GPUCA_GPUCODE void GPUTrackingRefitProcessor::InitializeProcessor() {} void GPUTrackingRefitProcessor::RegisterMemoryAllocation() @@ -55,6 +56,7 @@ void GPUTrackingRefitProcessor::RegisterMemoryAllocation() void GPUTrackingRefitProcessor::SetMaxData(const GPUTrackingInOutPointers& io) { } +#endif namespace { @@ -71,7 +73,7 @@ struct refitTrackTypes { } // anonymous namespace template <> -void GPUTrackingRefit::initProp(GPUTPCGMPropagator& prop) +GPUd() void GPUTrackingRefit::initProp(GPUTPCGMPropagator& prop) { prop.SetMaterialTPC(); prop.SetMaxSinPhi(GPUCA_MAX_SIN_PHI); @@ -80,20 +82,21 @@ void GPUTrackingRefit::initProp(GPUTPCGMPropagator& prop) prop.SetFitInProjections(mPparam->rec.fitInProjections != 0); prop.SetPropagateBzOnly(false); prop.SetPolynomialField(&mPparam->polynomialField); + prop.SetMatLUT(mPmatLUT); } template <> -void GPUTrackingRefit::initProp(const Propagator*& prop) +GPUd() void GPUTrackingRefit::initProp(const Propagator*& prop) { prop = mPpropagator; } template -void GPUTrackingRefit::convertTrack(T& trk, const S& trkX, U& prop, float* chi2) +GPUd() void GPUTrackingRefit::convertTrack(T& trk, const S& trkX, U& prop, float* chi2) { trk = trkX; } -static void convertTrackParam(GPUTPCGMTrackParam& trk, const TrackParCov& trkX) +GPUd() static void convertTrackParam(GPUTPCGMTrackParam& trk, const TrackParCov& trkX) { for (int i = 0; i < 5; i++) { trk.Par()[i] = trkX.getParams()[i]; @@ -103,7 +106,7 @@ static void convertTrackParam(GPUTPCGMTrackParam& trk, const TrackParCov& trkX) } trk.X() = trkX.getX(); } -static void convertTrackParam(TrackParCov& trk, const GPUTPCGMTrackParam& trkX) +GPUd() static void convertTrackParam(TrackParCov& trk, const GPUTPCGMTrackParam& trkX) { for (int i = 0; i < 5; i++) { trk.setParam(trkX.GetPar()[i], i); @@ -115,20 +118,20 @@ static void convertTrackParam(TrackParCov& trk, const GPUTPCGMTrackParam& trkX) } // Generic template <> -void GPUTrackingRefit::convertTrack(GPUTPCGMTrackParam& trk, const TrackParCov& trkX, GPUTPCGMPropagator& prop, float* chi2) +GPUd() void GPUTrackingRefit::convertTrack(GPUTPCGMTrackParam& trk, const TrackParCov& trkX, GPUTPCGMPropagator& prop, float* chi2) { convertTrackParam(trk, trkX); prop.SetTrack(&trk, trkX.getAlpha()); } template <> -void GPUTrackingRefit::convertTrack(TrackParCov& trk, const GPUTPCGMTrackParam& trkX, GPUTPCGMPropagator& prop, float* chi2) +GPUd() void GPUTrackingRefit::convertTrack(TrackParCov& trk, const GPUTPCGMTrackParam& trkX, GPUTPCGMPropagator& prop, float* chi2) { convertTrackParam(trk, trkX); trk.setAlpha(prop.GetAlpha()); } // GPUTPCGMMergedTrack input template <> -void GPUTrackingRefit::convertTrack(TrackParCov& trk, const GPUTPCGMMergedTrack& trkX, const Propagator*& prop, float* chi2) +GPUd() void GPUTrackingRefit::convertTrack(TrackParCov& trk, const GPUTPCGMMergedTrack& trkX, const Propagator*& prop, float* chi2) { initProp(prop); convertTrackParam(trk, trkX.GetParam()); @@ -136,62 +139,62 @@ void GPUTrackingRefit::convertTrack -void GPUTrackingRefit::convertTrack(GPUTPCGMMergedTrack& trk, const TrackParCov& trkX, const Propagator*& prop, float* chi2) +GPUd() void GPUTrackingRefit::convertTrack(GPUTPCGMMergedTrack& trk, const TrackParCov& trkX, const Propagator*& prop, float* chi2) { convertTrackParam(trk.Param(), trkX); trk.SetAlpha(trkX.getAlpha()); trk.Param().SetChi2(*chi2); } template <> -void GPUTrackingRefit::convertTrack(GPUTPCGMTrackParam& trk, const GPUTPCGMMergedTrack& trkX, GPUTPCGMPropagator& prop, float* chi2) +GPUd() void GPUTrackingRefit::convertTrack(GPUTPCGMTrackParam& trk, const GPUTPCGMMergedTrack& trkX, GPUTPCGMPropagator& prop, float* chi2) { initProp(prop); trk = trkX.GetParam(); prop.SetTrack(&trk, trkX.GetAlpha()); } template <> -void GPUTrackingRefit::convertTrack(GPUTPCGMMergedTrack& trk, const GPUTPCGMTrackParam& trkX, GPUTPCGMPropagator& prop, float* chi2) +GPUd() void GPUTrackingRefit::convertTrack(GPUTPCGMMergedTrack& trk, const GPUTPCGMTrackParam& trkX, GPUTPCGMPropagator& prop, float* chi2) { trk.SetParam(trkX); trk.SetAlpha(prop.GetAlpha()); } // TrackTPC input template <> -void GPUTrackingRefit::convertTrack(TrackParCov& trk, const TrackTPC& trkX, const Propagator*& prop, float* chi2) +GPUd() void GPUTrackingRefit::convertTrack(TrackParCov& trk, const TrackTPC& trkX, const Propagator*& prop, float* chi2) { initProp(prop); convertTrack(trk, trkX, prop, nullptr); *chi2 = trkX.getChi2(); } template <> -void GPUTrackingRefit::convertTrack(TrackTPC& trk, const TrackParCov& trkX, const Propagator*& prop, float* chi2) +GPUd() void GPUTrackingRefit::convertTrack(TrackTPC& trk, const TrackParCov& trkX, const Propagator*& prop, float* chi2) { convertTrack(trk, trkX, prop, nullptr); trk.setChi2(*chi2); } template <> -void GPUTrackingRefit::convertTrack(GPUTPCGMTrackParam& trk, const TrackTPC& trkX, GPUTPCGMPropagator& prop, float* chi2) +GPUd() void GPUTrackingRefit::convertTrack(GPUTPCGMTrackParam& trk, const TrackTPC& trkX, GPUTPCGMPropagator& prop, float* chi2) { initProp(prop); convertTrack(trk, trkX, prop, nullptr); trk.SetChi2(trkX.getChi2()); } template <> -void GPUTrackingRefit::convertTrack(TrackTPC& trk, const GPUTPCGMTrackParam& trkX, GPUTPCGMPropagator& prop, float* chi2) +GPUd() void GPUTrackingRefit::convertTrack(TrackTPC& trk, const GPUTPCGMTrackParam& trkX, GPUTPCGMPropagator& prop, float* chi2) { convertTrack(trk, trkX, prop, nullptr); trk.setChi2(trkX.GetChi2()); } // TrackParCovWithArgs input template <> -void GPUTrackingRefit::convertTrack(TrackParCov& trk, const GPUTrackingRefit::TrackParCovWithArgs& trkX, const Propagator*& prop, float* chi2) +GPUd() void GPUTrackingRefit::convertTrack(TrackParCov& trk, const GPUTrackingRefit::TrackParCovWithArgs& trkX, const Propagator*& prop, float* chi2) { initProp(prop); convertTrack(trk, trkX.trk, prop, nullptr); *chi2 = trkX.chi2 ? *trkX.chi2 : 0.f; } template <> -void GPUTrackingRefit::convertTrack(GPUTrackingRefit::TrackParCovWithArgs& trk, const TrackParCov& trkX, const Propagator*& prop, float* chi2) +GPUd() void GPUTrackingRefit::convertTrack(GPUTrackingRefit::TrackParCovWithArgs& trk, const TrackParCov& trkX, const Propagator*& prop, float* chi2) { convertTrack(trk.trk, trkX, prop, nullptr); if (trk.chi2) { @@ -199,14 +202,14 @@ void GPUTrackingRefit::convertTrack -void GPUTrackingRefit::convertTrack(GPUTPCGMTrackParam& trk, const GPUTrackingRefit::TrackParCovWithArgs& trkX, GPUTPCGMPropagator& prop, float* chi2) +GPUd() void GPUTrackingRefit::convertTrack(GPUTPCGMTrackParam& trk, const GPUTrackingRefit::TrackParCovWithArgs& trkX, GPUTPCGMPropagator& prop, float* chi2) { initProp(prop); convertTrack(trk, trkX.trk, prop, nullptr); trk.SetChi2(trkX.chi2 ? *trkX.chi2 : 0.f); } template <> -void GPUTrackingRefit::convertTrack(GPUTrackingRefit::TrackParCovWithArgs& trk, const GPUTPCGMTrackParam& trkX, GPUTPCGMPropagator& prop, float* chi2) +GPUd() void GPUTrackingRefit::convertTrack(GPUTrackingRefit::TrackParCovWithArgs& trk, const GPUTPCGMTrackParam& trkX, GPUTPCGMPropagator& prop, float* chi2) { convertTrack(trk.trk, trkX, prop, chi2); if (trk.chi2) { @@ -214,12 +217,13 @@ void GPUTrackingRefit::convertTrack GPUd() int GPUTrackingRefit::RefitTrack(T& trkX, bool outward, bool resetCov) { +#ifndef __OPENCL__ CADEBUG(int ii; printf("\nRefitting track\n")); typename refitTrackTypes::propagator prop; S trk; @@ -354,8 +358,8 @@ GPUd() int GPUTrackingRefit::RefitTrack(T& trkX, bool outward, bool resetCov) TrackParCovChi2 = 0.f; } CADEBUG(printf("\t%21sPropaga Alpha %8.3f , X %8.3f - Y %8.3f, Z %8.3f - QPt %7.2f (%7.2f), SP %5.2f (%5.2f) --- Res %8.3f %8.3f --- Cov sY %8.3f sZ %8.3f sSP %8.3f sPt %8.3f - YPt %8.3f\n", "", trk.getAlpha(), x, trk.getParams()[0], trk.getParams()[1], trk.getParams()[4], trk.getParams()[4], trk.getParams()[2], trk.getParams()[2], trk.getParams()[0] - y, trk.getParams()[1] - z, sqrtf(trk.getCov()[0]), sqrtf(trk.getCov()[2]), sqrtf(trk.getCov()[5]), sqrtf(trk.getCov()[14]), trk.getCov()[10])); - std::array p = {y, z}; - std::array c = {0, 0, 0}; + gpu::gpustd::array p = {y, z}; + gpu::gpustd::array c = {0, 0, 0}; mPparam->GetClusterErrors2(currentRow, z, getPar(trk)[2], getPar(trk)[3], c[0], c[2]); mPparam->UpdateClusterError2ByState(clusterState, c[0], c[2]); TrackParCovChi2 += trk.getPredictedChi2(p, c); @@ -376,6 +380,8 @@ GPUd() int GPUTrackingRefit::RefitTrack(T& trkX, bool outward, bool resetCov) trk.NormalizeAlpha(alpha); prop.SetAlpha(alpha); } else if constexpr (std::is_same::value) { + static constexpr float kDeg2Rad = M_PI / 180.f; + static constexpr float kSectAngle = 2 * M_PI / 18.f; if (mPparam->rec.TrackReferenceX <= 500) { if (prop->PropagateToXBxByBz(trk, mPparam->rec.TrackReferenceX)) { if (CAMath::Abs(trk.getY()) > trk.getX() * CAMath::Tan(kSectAngle / 2.f)) { @@ -391,6 +397,9 @@ GPUd() int GPUTrackingRefit::RefitTrack(T& trkX, bool outward, bool resetCov) convertTrack(trkX, trk, prop, &TrackParCovChi2); return nFitted; +#else + return 0; // TODO: Fixme, implement std::isSame for opencl +#endif } template GPUd() int GPUTrackingRefit::RefitTrack(GPUTPCGMMergedTrack& trk, bool outward, bool resetCov); @@ -400,20 +409,19 @@ template GPUd() int GPUTrackingRefit::RefitTrack(T template GPUd() int GPUTrackingRefit::RefitTrack(GPUTrackingRefit::TrackParCovWithArgs& trk, bool outward, bool resetCov); template GPUd() int GPUTrackingRefit::RefitTrack(GPUTrackingRefit::TrackParCovWithArgs& trk, bool outward, bool resetCov); -void GPUTrackingRefit::SetPtrsFromGPUConstantMem(const GPUConstantMem* v) +#ifndef GPUCA_GPUCODE +void GPUTrackingRefit::SetPtrsFromGPUConstantMem(const GPUConstantMem* v, MEM_CONSTANT(GPUParam) * p) { mPclusterState = v->ioPtrs.mergedTrackHitStates; mPclusterNative = v->ioPtrs.clustersNative; mPtrackHits = v->ioPtrs.mergedTrackHits; mPfastTransform = v->calibObjects.fastTransform; - mPparam = &v->param; + mPmatLUT = v->calibObjects.matLUT; + mPparam = p ? p : &v->param; } void GPUTrackingRefit::SetPropagatorDefault() { -#ifndef GPUCA_STANDALONE - mPpropagator = Propagator::Instance(); -#else - throw std::runtime_error("unsupported"); -#endif + mPpropagator = mPparam->GetDefaultO2Propagator(false); } +#endif diff --git a/GPU/GPUTracking/Refit/GPUTrackingRefit.h b/GPU/GPUTracking/Refit/GPUTrackingRefit.h index 7dbc710beeecc..0a8970dabaf6c 100644 --- a/GPU/GPUTracking/Refit/GPUTrackingRefit.h +++ b/GPU/GPUTracking/Refit/GPUTrackingRefit.h @@ -16,7 +16,6 @@ #include "GPUDef.h" #include "GPUProcessor.h" -#include namespace o2::dataformats { @@ -32,6 +31,7 @@ using TrackParCov = TrackParametrizationWithError; namespace o2::base { class Propagator; +class MatLayerCylSet; } // namespace o2::base namespace o2::tpc { @@ -55,7 +55,7 @@ class GPUTrackingRefit { public: void SetClusterStateArray(const unsigned char* v) { mPclusterState = v; } - void SetPtrsFromGPUConstantMem(const GPUConstantMem* v); + void SetPtrsFromGPUConstantMem(const GPUConstantMem* v, MEM_CONSTANT(GPUParam) * p = nullptr); void SetPropagator(const o2::base::Propagator* v) { mPpropagator = v; } void SetPropagatorDefault(); void SetClusterNative(const o2::tpc::ClusterNativeAccess* v) { mPclusterNative = v; } @@ -90,6 +90,7 @@ class GPUTrackingRefit private: const unsigned char* mPclusterState = nullptr; // Ptr to shared cluster state const o2::base::Propagator* mPpropagator = nullptr; // Ptr to propagator for TrackParCov track model + const o2::base::MatLayerCylSet* mPmatLUT = nullptr; // Ptr to material LUT const o2::tpc::ClusterNativeAccess* mPclusterNative = nullptr; // Ptr to cluster native access structure const GPUTPCGMMergedTrackHit* mPtrackHits = nullptr; // Ptr to hits for GPUTPCGMMergedTrack tracks const unsigned int* mPtrackHitReferences = nullptr; // Ptr to hits for TrackTPC tracks @@ -98,9 +99,9 @@ class GPUTrackingRefit template GPUd() int RefitTrack(T& trk, bool outward, bool resetCov); template - void convertTrack(T& trk, const S& trkX, U& prop, float* chi2); + GPUd() void convertTrack(T& trk, const S& trkX, U& prop, float* chi2); template - void initProp(U& prop); + GPUd() void initProp(U& prop); }; class GPUTrackingRefitProcessor : public GPUTrackingRefit, public GPUProcessor @@ -111,6 +112,7 @@ class GPUTrackingRefitProcessor : public GPUTrackingRefit, public GPUProcessor void RegisterMemoryAllocation(); void SetMaxData(const GPUTrackingInOutPointers& io); #endif + GPUTPCGMMergedTrack* mPTracks = nullptr; }; } // namespace o2::gpu diff --git a/GPU/GPUTracking/Refit/GPUTrackingRefitKernel.cxx b/GPU/GPUTracking/Refit/GPUTrackingRefitKernel.cxx new file mode 100644 index 0000000000000..de4f94aa586c3 --- /dev/null +++ b/GPU/GPUTracking/Refit/GPUTrackingRefitKernel.cxx @@ -0,0 +1,41 @@ +// 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 GPUTrackingRefitKernel.cxx +/// \author David Rohr + +#include "GPUTrackingRefitKernel.h" +#include "GPUTrackingRefit.h" + +using namespace GPUCA_NAMESPACE::gpu; + +template +GPUdii() void GPUTrackingRefitKernel::Thread(int nBlocks, int nThreads, int iBlock, int iThread, GPUsharedref() GPUSharedMemory& GPUrestrict() smem, processorType& GPUrestrict() processors) +{ + auto& refit = processors.trackingRefit; + for (unsigned int i = get_global_id(0); i < processors.ioPtrs.nMergedTracks; i += get_global_size(0)) { + if (refit.mPTracks[i].OK()) { + GPUTPCGMMergedTrack trk = refit.mPTracks[i]; + int retval; + if constexpr (I == mode0asGPU) { + retval = refit.RefitTrackAsGPU(trk, false, true); + } else if constexpr (I == mode1asTrackParCov) { + retval = refit.RefitTrackAsTrackParCov(trk, false, true); + } + if (retval > 0) { + refit.mPTracks[i] = trk; + } else { + refit.mPTracks[i].SetOK(false); + } + } + } +} +template GPUd() void GPUTrackingRefitKernel::Thread<0>(int nBlocks, int nThreads, int iBlock, int iThread, GPUsharedref() GPUSharedMemory& GPUrestrict() smem, processorType& GPUrestrict() processors); +template GPUd() void GPUTrackingRefitKernel::Thread<1>(int nBlocks, int nThreads, int iBlock, int iThread, GPUsharedref() GPUSharedMemory& GPUrestrict() smem, processorType& GPUrestrict() processors); diff --git a/GPU/GPUTracking/Refit/GPUTrackingRefitKernel.h b/GPU/GPUTracking/Refit/GPUTrackingRefitKernel.h new file mode 100644 index 0000000000000..4b2e5fdc16b6c --- /dev/null +++ b/GPU/GPUTracking/Refit/GPUTrackingRefitKernel.h @@ -0,0 +1,39 @@ +// 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 GPUTrackingRefitKernel.h +/// \author David Rohr + +#ifndef GPUTRACKINGREFITKERNEL_H +#define GPUTRACKINGREFITKERNEL_H + +#include "GPUGeneralKernels.h" +#include "GPUConstantMem.h" + +namespace o2::gpu +{ + +class GPUTrackingRefitKernel : public GPUKernelTemplate +{ + public: + GPUhdi() CONSTEXPR static GPUDataTypes::RecoStep GetRecoStep() { return GPUDataTypes::RecoStep::TPCCompression; } + + enum K : int { + mode0asGPU = 0, + mode1asTrackParCov = 1, + }; + + template + GPUd() static void Thread(int nBlocks, int nThreads, int iBlock, int iThread, GPUsharedref() GPUSharedMemory& GPUrestrict() smem, processorType& GPUrestrict() processors); +}; + +} // namespace o2::gpu + +#endif diff --git a/GPU/GPUTracking/Standalone/standalone.cxx b/GPU/GPUTracking/Standalone/standalone.cxx index d589ab09a803c..d859367c2490e 100644 --- a/GPU/GPUTracking/Standalone/standalone.cxx +++ b/GPU/GPUTracking/Standalone/standalone.cxx @@ -445,6 +445,16 @@ int SetupReconstruction() } } +#ifdef HAVE_O2HEADERS + chainTracking->SetDefaultO2PropagatorForGPU(); + if (configStandalone.testSyncAsync) { + chainTrackingAsync->SetDefaultO2PropagatorForGPU(); + } + if (configStandalone.proc.doublePipeline) { + chainTrackingPipeline->SetDefaultO2PropagatorForGPU(); + } +#endif + if (rec->Init()) { printf("Error initializing GPUReconstruction!\n"); return 1; diff --git a/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFDecodeZS.cxx b/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFDecodeZS.cxx index 1b17d0a8d0ffa..ae9d5ab3fda0c 100644 --- a/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFDecodeZS.cxx +++ b/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFDecodeZS.cxx @@ -170,7 +170,7 @@ GPUdii() void GPUTPCCFDecodeZS::decode(GPUTPCClusterFinder& clusterer, GPUShared positions[nDigitsTmp++] = pos; if (inFragment) { float q = float(byte & mask) * decodeBitsFactor; - q *= clusterer.getGainCorrection(row, pad); + q *= clusterer.GetConstantMem()->calibObjects.tpcPadGain->getGainCorrection(slice, row, pad); chargeMap[pos] = PackedCharge(q); } pad++; diff --git a/GPU/GPUTracking/TPCClusterFinder/GPUTPCClusterFinder.h b/GPU/GPUTracking/TPCClusterFinder/GPUTPCClusterFinder.h index 944c2be4de71c..d48f1e5d453c4 100644 --- a/GPU/GPUTracking/TPCClusterFinder/GPUTPCClusterFinder.h +++ b/GPU/GPUTracking/TPCClusterFinder/GPUTPCClusterFinder.h @@ -138,7 +138,6 @@ class GPUTPCClusterFinder : public GPUProcessor short mZSOffsetId = -1; short mOutputId = -1; - GPUdi() float getGainCorrection(tpccf::Row, tpccf::Pad) const; GPUdi() const GPUTPCGeometry* getGeometry() const; #ifndef GPUCA_GPUCODE