diff --git a/GPU/Common/GPUCommonMath.h b/GPU/Common/GPUCommonMath.h index f32cf61a057ae..2f8b0a3a25a28 100644 --- a/GPU/Common/GPUCommonMath.h +++ b/GPU/Common/GPUCommonMath.h @@ -106,6 +106,7 @@ class GPUCommonMath GPUd() constexpr static double Copysignd(double x, double y); GPUd() constexpr static float TwoPi() { return 6.2831853f; } GPUd() constexpr static float Pi() { return 3.1415927f; } + GPUd() constexpr static float Deg2Rad() { return Pi() / 180.f; } GPUd() constexpr static float Round(float x); GPUd() constexpr static float Floor(float x) { return GPUCA_CHOICE(floorf(x), floorf(x), floor(x)); } GPUd() static uint32_t Float2UIntReint(const float& x); diff --git a/GPU/GPUTracking/Base/GPUParam.h b/GPU/GPUTracking/Base/GPUParam.h index 30cc4d1e7dca9..5fab4b5e0fb54 100644 --- a/GPU/GPUTracking/Base/GPUParam.h +++ b/GPU/GPUTracking/Base/GPUParam.h @@ -80,8 +80,8 @@ struct GPUParam : public internal::GPUParam_t void UpdateRun3ClusterErrors(const float* yErrorParam, const float* zErrorParam); #endif - GPUd() constexpr uint32_t tpcMinHitsB5(float qPtB5) const { return CAMath::Abs(qPtB5) > 10 ? 10 : (CAMath::Abs(qPtB5) > 5 ? 15 : 29); } // Minimum hits should depend on Pt, low Pt tracks can have few hits. 29 Hits default, 15 for < 200 mev, 10 for < 100 mev - GPUd() constexpr float Alpha(int32_t iSector) const // TODO: Check if this is faster, or the lookup from GPUTPCGeometry + GPUd() constexpr uint32_t tpcMinHitsB5Abs(float qPtB5) const { return qPtB5 > 10.f ? 10 : (qPtB5 > 5.f ? (uint32_t)(20 - qPtB5) : (qPtB5 > 3.6f ? (uint32_t)(65 - 10.f * qPtB5) : 29)); } + GPUd() constexpr float Alpha(int32_t iSector) const // TODO: Check if this is faster, or the lookup from GPUTPCGeometry { if (iSector >= (int32_t)o2::tpc::constants::MAXSECTOR / 2) { iSector -= o2::tpc::constants::MAXSECTOR / 2; diff --git a/GPU/GPUTracking/Base/GPUReconstruction.cxx b/GPU/GPUTracking/Base/GPUReconstruction.cxx index 3e7509d8fad73..c1a09f4075bc9 100644 --- a/GPU/GPUTracking/Base/GPUReconstruction.cxx +++ b/GPU/GPUTracking/Base/GPUReconstruction.cxx @@ -295,7 +295,7 @@ int32_t GPUReconstruction::InitPhaseBeforeDevice() if (GetProcessingSettings().deterministicGPUReconstruction && GetProcessingSettings().debugLevel >= 6) { mProcessingSettings->nTPCClustererLanes = 1; } - if (GetProcessingSettings().createO2Output > 1 && GetProcessingSettings().runQA && GetProcessingSettings().qcRunFraction == 100.f) { + if (GetProcessingSettings().createO2Output > 1 && ((GetProcessingSettings().runQA && GetProcessingSettings().qcRunFraction == 100.f) || GetProcessingSettings().debug.mergerDumpTrackNCl)) { mProcessingSettings->createO2Output = 1; } if (!GetProcessingSettings().createO2Output || !IsGPU()) { diff --git a/GPU/GPUTracking/Base/GPUReconstructionCPU.cxx b/GPU/GPUTracking/Base/GPUReconstructionCPU.cxx index 9ee6fae1e0fd9..a6d7af2d812aa 100644 --- a/GPU/GPUTracking/Base/GPUReconstructionCPU.cxx +++ b/GPU/GPUTracking/Base/GPUReconstructionCPU.cxx @@ -216,10 +216,10 @@ int32_t GPUReconstructionCPU::ExitDevice() int32_t GPUReconstructionCPU::RunChains() { mMemoryScalers->temporaryFactor = 1.; - if (GetProcessingSettings().memoryScalingFuzz) { + if (GetProcessingSettings().debug.memoryScalingFuzz) { static std::mt19937 rng; static std::uniform_int_distribution dist(0, 1000000); - uint64_t fuzzFactor = GetProcessingSettings().memoryScalingFuzz == 1 ? dist(rng) : GetProcessingSettings().memoryScalingFuzz; + uint64_t fuzzFactor = GetProcessingSettings().debug.memoryScalingFuzz == 1 ? dist(rng) : GetProcessingSettings().debug.memoryScalingFuzz; GPUInfo("Fuzzing memory scaling factor with %lu", fuzzFactor); mMemoryScalers->fuzzScalingFactor(fuzzFactor); } diff --git a/GPU/GPUTracking/DataTypes/GPUTPCGeometry.h b/GPU/GPUTracking/DataTypes/GPUTPCGeometry.h index ed334aef3a75f..ac2560a45ef90 100644 --- a/GPU/GPUTracking/DataTypes/GPUTPCGeometry.h +++ b/GPU/GPUTracking/DataTypes/GPUTPCGeometry.h @@ -16,7 +16,11 @@ #define GPUTPCGEOMETRY_H #include "GPUCommonDef.h" +#include "GPUCommonMath.h" #include "DataFormatsTPC/Constants.h" +#ifndef GPUCA_GPUCODE_DEVICE +#include +#endif namespace o2::gpu { @@ -174,6 +178,8 @@ class GPUTPCGeometry const float v = (sector >= NSECTORS / 2) ? -z : z; return (250.f - v) * FACTOR_Z2T; // Used in compression, must remain constant at 250cm } + + GPUd() static constexpr float kSectAngle() { return 2 * CAMath::Pi() / 18.f; } }; } // namespace o2::gpu diff --git a/GPU/GPUTracking/Definitions/GPUDef.h b/GPU/GPUTracking/Definitions/GPUDef.h index 54c180d5f7187..d3aa2acfa51d5 100644 --- a/GPU/GPUTracking/Definitions/GPUDef.h +++ b/GPU/GPUTracking/Definitions/GPUDef.h @@ -70,10 +70,23 @@ #ifdef CADEBUG #undef CADEBUG #endif - #if GPUCA_CADEBUG == 1 && !defined(GPUCA_GPUCODE) - #define CADEBUG(...) __VA_ARGS__ - #define CADEBUG2(cmd, ...) {__VA_ARGS__; cmd;} - #define GPUCA_CADEBUG_ENABLED + #ifdef CADEBUG2 + #undef CADEBUG2 + #endif + #ifdef CADEBUG3 + #undef CADEBUG3 + #endif + #if GPUCA_CADEBUG >= 1 && !defined(GPUCA_GPUCODE) + #if GPUCA_CADEBUG >= 2 + #define CADEBUG(...) if (CADEBUG_CONDITION) {__VA_ARGS__;} + #define CADEBUG2(cmd, ...) if (CADEBUG_CONDITION) {__VA_ARGS__; cmd;} + #define GPUCA_CADEBUG_ENABLED 2 + #else + #define CADEBUG(...) __VA_ARGS__ + #define CADEBUG2(cmd, ...) {__VA_ARGS__; cmd;} + #define GPUCA_CADEBUG_ENABLED 1 + #endif + #define CADEBUG3(...) __VA_ARGS__ #endif #undef GPUCA_CADEBUG #endif @@ -81,5 +94,6 @@ #ifndef CADEBUG #define CADEBUG(...) #define CADEBUG2(cmd, ...) {cmd;} + #define CADEBUG3(...) #endif // clang-format on diff --git a/GPU/GPUTracking/Definitions/GPUSettingsList.h b/GPU/GPUTracking/Definitions/GPUSettingsList.h index d1e13efc93518..3e053313136f3 100644 --- a/GPU/GPUTracking/Definitions/GPUSettingsList.h +++ b/GPU/GPUTracking/Definitions/GPUSettingsList.h @@ -32,7 +32,7 @@ BeginNamespace(o2::gpu) // Reconstruction parameters for TPC, no bool in here !!! BeginSubConfig(GPUSettingsRecTPC, tpc, configStandalone.rec, "RECTPC", 0, "Reconstruction settings", rec_tpc) -AddOptionRTC(rejectQPtB5, float, 1.f / 0.050f, "", 0, "QPt threshold to reject clusters of TPC tracks (Inverse Pt, scaled to B=0.5T!!!)") +AddOptionRTC(rejectQPtB5, float, 1.f / 0.050f, "", 0, "QPt threshold to reject clusters of TPC tracks (Inverse Pt, scaled to B=0.5T!!!)") // TODO: Sort these options automatically for parameter size AddOptionRTC(hitPickUpFactor, float, 1.f, "", 0, "multiplier for the combined cluster+track error during track following") AddOptionRTC(hitSearchArea2, float, 2.f, "", 0, "square of maximum search road of hits during seeding") AddOptionRTC(neighboursSearchArea, float, 3.f, "", 0, "area in cm for the search of neighbours, for z only used if searchWindowDZDR = 0") @@ -102,6 +102,20 @@ AddOptionRTC(rejectEdgeClustersMargin, float, 0.f, "", 0, "Margin in cm of Y pos AddOptionRTC(rejectEdgeClustersSigmaMargin, float, 0.f, "", 0, "Margin factor for trackSigmaY when rejecting edge clusters based on uncorrected track Y") AddOptionRTC(trackletMaxSharedFraction, float, 0.1f, "", 0, "Max fraction of shared clusters for tracklet") AddOptionRTC(trackletMinSharedNormFactor, float, 0.f, "", 0, "Max shared defined as trackletMinSharedNormFactor*max(current_nhits,trackletMinSharedNormFactor*minHits,1)") +AddOptionRTC(trackFitMinHitsMaxQPtIncrease, float, 0.2f, "", 0, "For the QPt-dependent minimum number of hits of the track fit, abs(QPt) is increased by 1 sigma, but at most by this fraction (0 = use fitted QPt)") +AddOptionRTC(rebuildTrackMaxSharedFraction, float, 0.1f, "", 0, "Max fraction of shared clusters for rebuilt tracks") +AddOptionRTC(rebuildTrackMaxNonIntCov, float, 4.f, "", 0, "Max Err2 allowed for non-interpolated cluster attachment during rebuild") +AddOptionRTC(rebuildErrorWeightFactor, float, 5.f, "", 0, "Factor for the error (distance^2) of a cluster candidate in its hit weight 1 / (1 + factor * error) during rebuild") +AddOptionRTC(maxSinPhiForFitUpdate, float, 0.9396926f, "", 0, "Max abs(sinPhi) of the track to update it with a cluster in the track fit)") // default sin(70 deg = 0.9396926 +AddOptionRTC(rebuildTrackMaxQPtDrift, float, 0.5f, "", 0, "Stop the track extrapolation during rebuild when QPt changed by more than this fraction of max(1, abs(QPt)) compared to before the extrapolation (0 = disabled)") +AddOptionRTC(rebuildTrackExtrMaxAvgChi2, float, 4.f, "", 0, "Max average chi2 per cluster to accept an extrapolated segment (0 = disabled)") +AddOptionRTC(rebuildTrackFallbackErr, float, 1.f, "", 0, "Error (distance^2) assigned to the original clusters of a track used as fallback candidates in rows that could not be searched during rebuild") +AddOptionRTC(rebuildTrackFallbackTrackChi2, float, 4.f, "", 0, "Chi2 per cluster assumed for the hit weight of tracks without successful fit (NDF <= 0) during rebuild") +AddOptionRTC(rebuildTrackAltCandidateMaxErrFactor, float, 2.f, "", 0, "Alternative cluster candidates of a row are used only if their distance^2 is below this factor times the distance^2 of the best candidate") +AddOptionRTC(rebuildTrackAltCandidateGoodErr, float, 0.25f, "", 0, "Alternative cluster candidates of a row with distance^2 below this value are always used") +AddOptionRTC(rebuildTrackErrFactorOneSided, float, 1.5f, "", 0, "Factor applied to the error (distance^2) of cluster candidates found with the track position from only one side during rebuild") +AddOptionRTC(rebuildTrackErrFactorExtrapolation, float, 1.f, "", 0, "Factor applied to the error (distance^2) of cluster candidates found by track extrapolation during rebuild") +AddOptionRTC(rebuildTrackMaxInterpolationDistY, float, 4.f, "", 0, "Max distance in Y between the track positions from both sides to use their interpolation for the cluster search in rows without cluster during rebuild") AddOptionRTC(maxTimeBinAboveThresholdIn1000Bin, uint16_t, 500, "", 0, "Except pad from cluster finding if total number of charges in a fragment is above this baseline (disable = 0)") AddOptionRTC(maxConsecTimeBinAboveThreshold, uint16_t, 200, "", 0, "Except pad from cluster finding if number of consecutive charges in a fragment is above this baseline (disable = 0)") AddOptionRTC(noisyPadSaturationThreshold, uint16_t, 700, "", 0, "Threshold where a timebin is considered saturated, disabling the noisy pad check for that pad") @@ -135,12 +149,16 @@ AddOptionRTC(cfNoiseSuppressionEpsilon, uint8_t, 10, "", 0, "Cluster Finder: Dif AddOptionRTC(cfNoiseSuppressionEpsilonRelative, uint8_t, 76, "", 0, "Cluster Finder: Difference between peak and charge for the charge to count as a minima during noise suppression, relative as fraction of 255") AddOptionRTC(cfEdgeTwoPads, uint8_t, 0, "", 0, "Flag clusters with peak on the 2 pads closes to the sector edge as edge cluster") AddOptionRTC(nWays, uint8_t, 3, "", 0, "Do N fit passes in final fit of merger (must be odd to end with inward fit)") -AddOptionRTC(trackFitRejectMode, int8_t, 5, "", 0, "0: no limit on rejection or missed hits, >0: break after n rejected hits, <0: reject at max -n hits") +AddOptionRTC(rebuildTrackInFit, uint8_t, 1, "", 0, "Rebuild track completely during fit based on clusters closed to interpolated track positions") +AddOptionRTC(rebuildTrackInFitClusterCandidates, uint8_t, 3, "", 0, "Number of cluster candidates per row for rebuilt track") +AddOptionRTC(disableRebuildAttachment, uint8_t, 0, "", 0, "Bitmask to disable certain attachment steps during track rebuid (1: current row, 2: interpolate missing row, 4: one sided interpolation, 8: original hit, 16: extrapolation)") +AddOptionRTC(disableMarkAdjacent, uint8_t, 0, "", 0, "Bitmask to disable certain steps during refit to mark adjacenrt clusters (1: current row, 2: extrapolate missing rows, 4: loop following)") // TODO: Add history +AddOptionRTC(trackFitRejectMode, int8_t, 5, "", 0, "0: no limit on rejection or missed hits, >0: break after n rejected hits") AddOptionRTC(rejectIFCLowRadiusCluster, uint8_t, 1, "", 0, "Reject clusters that get the IFC mask error during refit") AddOptionRTC(dEdxTruncLow, uint8_t, 2, "", 0, "Low truncation threshold, fraction of 128") AddOptionRTC(dEdxTruncHigh, uint8_t, 77, "", 0, "High truncation threshold, fraction of 128") AddOptionRTC(extrapolationTracking, int8_t, 1, "", 0, "Enable Extrapolation Tracking (prolong tracks to adjacent sectors to find short segments)") -AddOptionRTC(disableRefitAttachment, uint8_t, 0, "", 0, "Bitmask to disable certain attachment steps during refit (1: attachment, 2: propagation, 4: loop following, 8: mirroring)") +AddOptionRTC(looperFollowMode, uint8_t, 1, "", 0, "0 = simple, 1 = advances, for disabling use disableMarkAdjacent = 4") AddOptionRTC(rejectionStrategy, uint8_t, o2::gpu::GPUSettings::RejectionStrategyA, "", 0, "Enable rejection of TPC clusters for compression (0 = no, 1 = strategy A, 2 = strategy B)") AddOptionRTC(mergeLoopersAfterburner, uint8_t, 1, "", 0, "Run afterburner for additional looper merging") AddOptionRTC(compressionTypeMask, uint8_t, o2::gpu::GPUSettings::CompressionFull, "", 0, "TPC Compression mode bits (1=truncate charge/width LSB, 2=differences, 4=track-model)") @@ -152,7 +170,7 @@ AddOptionRTC(mergerInterpolateErrors, uint8_t, 1, "", 0, "Use interpolation inst AddOptionRTC(mergerInterpolateRejectAlsoOnCurrentPosition, uint8_t, 1, "", 0, "When using mergerInterpolateErrors, reject based on chi2 twice computed with interpolated and current track position starting from NDF > mergerNonInterpolateRejectMinNDF") AddOptionRTC(mergerNonInterpolateRejectMinNDF, uint8_t, 5, "", 0, "Minimum NDF of track for non-interpolated reject (both for chi2 and absolute distance)") AddOptionRTC(mergeCE, uint8_t, 1, "", 0, "Merge tracks accross the central electrode") -AddOptionRTC(retryRefit, int8_t, 1, "", 0, "Retry refit with seeding errors and without cluster rejection when fit fails (=2 means retry in same kernel, =1 for separate kernel") +AddOptionRTC(retryRefit, int8_t, 1, "", 0, "Retry refit with seeding errors and without cluster rejection when fit fails") AddOptionRTC(enablePID, int8_t, 1, "", 0, "Enable PID response") AddOptionRTC(PID_useNsigma, int8_t, 1, "", 0, "Use nSigma instead of absolute distance in PID response") AddOptionRTC(adddEdxSubThresholdClusters, int8_t, 1, "", 0, "Add sub threshold clusters in TPC dEdx computation") @@ -162,8 +180,16 @@ AddOptionRTC(rejectEdgeClustersInSeeding, int8_t, 0, "", 0, "Reject edge cluster AddOptionRTC(rejectEdgeClustersInTrackFit, int8_t, 0, "", 0, "Reject edge clusters based on uncorrected track Y during track fit") AddOptionRTC(tubeExtraProtectMinRow, uint8_t, 20, "", 0, "Increase Protection, decrease removal by factor 2, when below this row") AddOptionRTC(tubeExtraProtectEdgePads, uint8_t, 2, "", 0, "Increase Protection, decrease removal by factor 2, when on this number of pads from the edge") - -AddOptionArray(PID_remap, int8_t, 9, (0, 1, 2, 3, 4, 5, 6, 7, 8), "", 0, "Remap Ipid to PID_reamp[Ipid] (no remap if<0)") // BUG: CUDA cannot yet hand AddOptionArrayRTC +AddOptionRTC(rebuildTrackExtrMaxMissingRows, uint8_t, 20, "", 0, "Maximum total number of rows allowed to be missing during track extrapolation") +AddOptionRTC(rebuildTrackExtrMinConsecGoodRows, uint8_t, 8, "", 0, "Minimum number of consecutive rows required to accept extrapolated segment") +AddOptionRTC(rebuildTrackExtrMinConsecGoodRowsEdge, uint8_t, 3, "", 0, "Minimum number of consecutive rows required to accept extrapolated segment") +AddOptionRTC(rebuildTrackExtrMaxRowsBack, uint8_t, 3, "", 0, "After the track extrapolation, propagate the track back to the last accepted cluster if it is at most this number of rows away, otherwise revert to the track before the extrapolation") +AddOptionRTC(rebuildTrackExtrConsecGoodRowsMaxGap, uint8_t, 1, "", 0, "Max row gap size to consider the segment as consecutive rows") +AddOptionRTC(rebuildTrackCandidateErrChi2, uint8_t, 0, "", 0, "Rank cluster candidates during rebuild by chi2 instead of distance^2 (rebuildTrackAltCandidateGoodErr is then in units of chi2)") +AddOptionRTC(rebuildTrackMaxRowsFromEnd, uint8_t, 10, "", 0, "If the last fitted hit of the outward way of the rebuild is more than this number of rows away from the end of the track, the final fit starts from the original track parameters") +AddOptionRTC(rebuildTrackFallbackMaxNCl, uint8_t, 20, "", 0, "For tracks with at most this number of clusters, keep the original cluster of a row as candidate during rebuild if no candidate was found but it passes the interpolated chi2 cut (0 = disabled)") +AddOptionRTC(rebuildTrackNonIntMinNDF, uint8_t, 10, "", 0, "NDF the track must exceed for non-interpolated cluster attachment (one-sided search and extrapolation) during rebuild") +AddOptionArray(PID_remap, int8_t, 9, (0, 1, 2, 3, 4, 5, 6, 7, 8), "", 0, "Remap Ipid to PID_reamp[Ipid] (no remap if<0)") // BUG: CUDA cannot yet handle AddOptionArrayRTC AddHelp("help", 'h') EndConfig() @@ -322,6 +348,14 @@ AddOption(tpcDecodingSafetyBuffer, uint16_t, 1000, "", 0, "Scaling Factor: safet AddHelp("help", 'h') EndConfig() +// Debug Settings +BeginSubConfig(GPUSettingsProcessingDebug, debug, configStandalone.proc, "DEBUG", 0, "Debugging Settings", proc_debug) +AddOption(memoryScalingFuzz, uint64_t, 0, "", 0, "Fuzz the memoryScalingFactor (0 disable, 1 enable, >1 set seed", def(1)) +AddOption(mergerMCLabels, bool, false, "", 0, "Create MC labels for merged tracks before refit for debugging") +AddOption(mergerDumpTrackNCl, bool, false, "", 0, "Dump file with nCl per merged track fit") +AddHelp("help", 'h') +EndConfig() + // Settings steering the processing once the device was selected, only available on the host BeginSubConfig(GPUSettingsProcessing, proc, configStandalone, "PROC", 0, "Processing settings", proc) AddOption(deviceNum, int32_t, -1, "gpuDevice", 0, "Set GPU device to use (-1: automatic, -2: for round-robin usage in timeslice-pipeline)") @@ -351,7 +385,6 @@ AddOption(memoryAllocationStrategy, int8_t, 0, "", 0, "Memory Allocation Strateg AddOption(forceMemoryPoolSize, uint64_t, 1, "memSize", 0, "Force size of allocated GPU / page locked host memory (0 = disable, 1 / 2 = auto, other values set the size of the memory pool)", min(0ul)) AddOption(forceHostMemoryPoolSize, uint64_t, 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(memoryScalingFuzz, uint64_t, 0, "", 0, "Fuzz the memoryScalingFactor (0 disable, 1 enable, >1 set seed", def(1)) AddOption(tpcInputWithClusterRejection, uint8_t, 0, "", 0, "Indicate whether the TPC input is CTF data with cluster rejection, to tune buffer estimations") AddOption(forceMaxMemScalers, uint64_t, 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") @@ -398,7 +431,7 @@ AddOption(tpcUseOldCPUDecoding, bool, false, "", 0, "Enable old CPU-based TPC de AddOption(tpcApplyCFCutsAtDecoding, bool, false, "", 0, "Apply cluster cuts from clusterization during decoding of compressed clusters") AddOption(tpcApplyClusterFilterOnCPU, uint8_t, 0, "", 0, "Apply custom cluster filter of GPUTPCClusterFilter class, 0: off, 1: debug, 2: PbPb23") AddOption(tpcWriteClustersAfterRejection, bool, false, "", 0, "Apply TPC rejection strategy before writing clusters") -AddOption(oclPlatformNum, int32_t, -1, "", 0, "Platform to use, in case the backend provides multiple platforms (OpenCL only, -1 = auto-select, -2 query all platforms (also incompatible))") +AddOption(oclPlatformNum, int32_t, -1, "", 0, "Platform to use, in case the backend provides multiple platforms (OpenCL only, -1 = auto-select, -2 query all platforms (also incompatible))") // TODO: Create some backend-specific options AddOption(oclCompileFromSources, bool, false, "", 0, "Compile OpenCL binary from included source code instead of using included spirv code") AddOption(oclOverrideSourceBuildFlags, std::string, "", "", 0, "Override OCL build flags for compilation from source, put a space for empty options") AddOption(metalOverrideSourceBuildFlags, std::string, "", "", 0, "Override Metal build flags for compilation from source, put a space for empty options") @@ -422,6 +455,7 @@ AddSubConfig(GPUSettingsProcessingRTCtechnical, rtctech) AddSubConfig(GPUSettingsProcessingParam, param) AddSubConfig(GPUSettingsProcessingNNclusterizer, nn) AddSubConfig(GPUSettingsProcessingScaling, scaling) +AddSubConfig(GPUSettingsProcessingDebug, debug) AddHelp("help", 'h') EndConfig() #endif // !__OPENCL__ && !__METAL__ diff --git a/GPU/GPUTracking/Definitions/Parameters/GPUParameters.csv b/GPU/GPUTracking/Definitions/Parameters/GPUParameters.csv index 330c9a3fa7377..fb2038c954f64 100644 --- a/GPU/GPUTracking/Definitions/Parameters/GPUParameters.csv +++ b/GPU/GPUTracking/Definitions/Parameters/GPUParameters.csv @@ -52,6 +52,11 @@ GPUTPCGMMergerPrepareForFit_step2,256,,256,256,,,,,,256,256,256,256,,256,,256,25 GPUTPCGMMergerFinalize_step0,256,,,256,,,,,,256,,256,256,,256,,256,256, GPUTPCGMMergerFinalize_step1,256,,,256,,,,,,256,,256,256,,256,,256,256, GPUTPCGMMergerFinalize_step2,256,,,256,,,,,,256,,256,256,,256,,256,256, +GPUTPCGMMergerHitWeights_prepare,256,,,,,,,,,,,,,,,,,, +GPUTPCGMMergerHitWeights_compute,256,,,,,,,,,,,,,,,,,, +GPUTPCGMMergerHitWeights_resolve1,256,,,,,,,,,,,,,,,,,, +GPUTPCGMMergerHitWeights_resolve2,256,,,,,,,,,,,,,,,,,, +GPUTPCGMMergerHitWeights_resolveShared,256,,,,,,,,,,,,,,,,,, GPUTPCGMMergerMergeLoopers_step0,256,,,,,,,,,256,,256,256,,256,,256,256, GPUTPCGMMergerMergeLoopers_step1,256,,,,,,,,,256,,256,256,,256,,256,256, GPUTPCGMMergerMergeLoopers_step2,256,,,,,,,,,256,,256,256,,256,,256,256, diff --git a/GPU/GPUTracking/GPUTrackingLinkDef_O2_DataTypes.h b/GPU/GPUTracking/GPUTrackingLinkDef_O2_DataTypes.h index 21e495e3ca6c6..2628821bb3ca1 100644 --- a/GPU/GPUTracking/GPUTrackingLinkDef_O2_DataTypes.h +++ b/GPU/GPUTracking/GPUTrackingLinkDef_O2_DataTypes.h @@ -33,6 +33,7 @@ #pragma link C++ class o2::gpu::internal::GPUConfigurableParamGPUSettingsProcessingRTCtechnical + ; #pragma link C++ class o2::gpu::internal::GPUConfigurableParamGPUSettingsProcessingNNclusterizer + ; #pragma link C++ class o2::gpu::internal::GPUConfigurableParamGPUSettingsProcessingScaling + ; +#pragma link C++ class o2::gpu::internal::GPUConfigurableParamGPUSettingsProcessingDebug + ; #pragma link C++ class o2::gpu::internal::GPUConfigurableParamGPUSettingsDisplay + ; #pragma link C++ class o2::gpu::internal::GPUConfigurableParamGPUSettingsDisplayInternal + ; #pragma link C++ class o2::gpu::internal::GPUConfigurableParamGPUSettingsDisplayLight + ; diff --git a/GPU/GPUTracking/Global/GPUChainTracking.cxx b/GPU/GPUTracking/Global/GPUChainTracking.cxx index 8a22e2c03ea47..df74187a4f9d5 100644 --- a/GPU/GPUTracking/Global/GPUChainTracking.cxx +++ b/GPU/GPUTracking/Global/GPUChainTracking.cxx @@ -266,6 +266,14 @@ bool GPUChainTracking::ValidateSettings() GPUError("Cannot do error interpolation with NWays < 3!"); return false; } + if (param().rec.tpc.rebuildTrackInFit && !param().rec.tpc.mergerInterpolateErrors) { + GPUError("Need error interpolation to rebuild tracks during fit"); + return false; + } + if (param().rec.tpc.rebuildTrackInFit && param().rec.tpc.rebuildTrackInFitClusterCandidates == 0) { + GPUError("Need at least 1 cluster candidate to rebuild tracks during fit"); + return false; + } if (param().continuousMaxTimeBin > (int32_t)GPUSettings::TPC_MAX_TF_TIME_BIN) { GPUError("configured max time bin exceeds 256 orbits"); return false; @@ -1023,6 +1031,7 @@ void GPUChainTracking::ApplySyncSettings(GPUSettingsProcessing& proc, GPUSetting { if (syncMode) { rec.useMatLUT = false; + rec.tpc.rebuildTrackMaxNonIntCov = 0.f; // TODO: Check if this yields a performance benefit } if (proc.rtc.optSpecialCode == -1) { proc.rtc.optSpecialCode = syncMode; diff --git a/GPU/GPUTracking/Global/GPUChainTracking.h b/GPU/GPUTracking/Global/GPUChainTracking.h index 1507147bbafda..92eda50194c74 100644 --- a/GPU/GPUTracking/Global/GPUChainTracking.h +++ b/GPU/GPUTracking/Global/GPUChainTracking.h @@ -322,6 +322,9 @@ class GPUChainTracking : public GPUChain int32_t RunTRDTrackingInternal(); uint32_t StreamForSector(uint32_t sector) const; + // Debug + void DoDebugMergerNClDump(const GPUTPCGMMerger& merger); + std::mutex mMutexUpdateCalib; std::unique_ptr mPipelineFinalizationCtx; GPUChainTrackingFinalContext* mPipelineNotifyCtx = nullptr; diff --git a/GPU/GPUTracking/Global/GPUChainTrackingDebugAndProfiling.cxx b/GPU/GPUTracking/Global/GPUChainTrackingDebugAndProfiling.cxx index 7659d19693777..13369405ac519 100644 --- a/GPU/GPUTracking/Global/GPUChainTrackingDebugAndProfiling.cxx +++ b/GPU/GPUTracking/Global/GPUChainTrackingDebugAndProfiling.cxx @@ -187,8 +187,8 @@ void GPUChainTracking::PrintMemoryRelations() GPUInfo("MEMREL SectorTrackHits NCl %d NTrkH %d", processors()->tpcTrackers[i].NHitsTotal(), *processors()->tpcTrackers[i].NTrackHits()); } if (processors()->tpcMerger.Memory()) { - GPUInfo("MEMREL Tracks NCl %d NTrk %d", processors()->tpcMerger.NMaxClusters(), processors()->tpcMerger.NMergedTracks()); - GPUInfo("MEMREL TrackHitss NCl %d NTrkH %d", processors()->tpcMerger.NMaxClusters(), processors()->tpcMerger.NMergedTrackClusters()); + GPUInfo("MEMREL Tracks NCl %d NTrk %d", processors()->tpcMerger.NClusters(), processors()->tpcMerger.NMergedTracks()); + GPUInfo("MEMREL TrackHitss NCl %d NTrkH %d", processors()->tpcMerger.NClusters(), processors()->tpcMerger.NMergedTrackClusters()); } } @@ -217,7 +217,7 @@ void GPUChainTracking::PrintKernelDebugOutput() void GPUChainTracking::PrintOutputStat() { int32_t nTracks = 0, nAttachedClusters = 0, nAttachedClustersFitted = 0, nAdjacentClusters = 0; - uint32_t nCls = GetProcessingSettings().doublePipeline ? mIOPtrs.clustersNative->nClustersTotal : processors()->tpcMerger.NMaxClusters(); + uint32_t nCls = GetProcessingSettings().doublePipeline ? mIOPtrs.clustersNative->nClustersTotal : processors()->tpcMerger.NClusters(); if (GetProcessingSettings().createO2Output > 1) { nTracks = mIOPtrs.nOutputTracksTPCO2; nAttachedClusters = mIOPtrs.nMergedTrackHits; @@ -419,3 +419,26 @@ void GPUChainTracking::DoDebugRawDump() mRec->DumpSettings((dirName + "/").c_str()); DumpData((dirName + "/event.0.dump").c_str(), &ioPtrs); } + +void GPUChainTracking::DoDebugMergerNClDump(const GPUTPCGMMerger& merger) +{ + if (GetProcessingSettings().debug.mergerDumpTrackNCl) { +#ifndef GPUCA_DETERMINISTIC_MODE + GPUError("GPUCA_DETERMINISTIC_MODE not enabled, track dump will not be deterministic"); +#endif + if (!GetProcessingSettings().deterministicGPUReconstruction) { + GPUError("deterministicGPUReconstruction not enabled, track dump will not be deterministic"); + } + if (GetRecoStepsGPU() & RecoStep::TPCMerging) { + SynchronizeGPU(); + } + FILE* fout = fopen("merger_tracks_ncl.txt", "w+"); + if (fout == nullptr) { + throw std::runtime_error("Error Opening Debug File"); + } + for (int32_t i = 0; i < merger.NMergedTracks(); i++) { + fprintf(fout, "Track %7d NCl %5u\n", i, merger.MergedTracks()[i].OK() ? merger.MergedTracks()[i].NClustersFitted() : 0u); + } + fclose(fout); + } +} diff --git a/GPU/GPUTracking/Global/GPUChainTrackingMerger.cxx b/GPU/GPUTracking/Global/GPUChainTrackingMerger.cxx index 528a848acf389..38d02d99277fd 100644 --- a/GPU/GPUTracking/Global/GPUChainTrackingMerger.cxx +++ b/GPU/GPUTracking/Global/GPUChainTrackingMerger.cxx @@ -96,7 +96,6 @@ int32_t GPUChainTracking::RunTPCTrackingMerger(bool synchronizeOutput) uint32_t numBlocks = (!mRec->IsGPU() || doGPU) ? BlockCount() : 1; GPUTPCGMMerger& Merger = processors()->tpcMerger; GPUTPCGMMerger& MergerShadow = doGPU ? processorsShadow()->tpcMerger : Merger; - GPUTPCGMMerger& MergerShadowAll = doGPU ? processorsShadow()->tpcMerger : Merger; const int32_t outputStream = OutputStream(); if (GetProcessingSettings().debugLevel >= 2) { GPUInfo("Running TPC Merger"); @@ -142,7 +141,7 @@ int32_t GPUChainTracking::RunTPCTrackingMerger(bool synchronizeOutput) // Merge within Sectors runKernel(GetGridAuto(0, deviceType), false); - runKernel({{1, -WarpSize(), 0, deviceType, RecoStep::TPCMerging}}, MergerShadowAll.TmpCounter(), NSECTORS * sizeof(*MergerShadowAll.TmpCounter())); + runKernel({{1, -WarpSize(), 0, deviceType, RecoStep::TPCMerging}}, MergerShadow.TmpCounter(), NSECTORS * sizeof(*MergerShadow.TmpCounter())); runKernel(GetGridAuto(0, deviceType)); RunTPCTrackingMerger_MergeBorderTracks(GPUTPCGMMerger::mergeModes::mergeWithinSector, deviceType); RunTPCTrackingMerger_Resolve(0, 1, deviceType); @@ -150,23 +149,23 @@ int32_t GPUChainTracking::RunTPCTrackingMerger(bool synchronizeOutput) // Merge between sectors - transport to the middle of the sector and rotate vertically to the border on the left / right runKernel(GetGridAuto(0, deviceType), false); - runKernel({{1, -WarpSize(), 0, deviceType, RecoStep::TPCMerging}}, MergerShadowAll.TmpCounter(), 2 * NSECTORS * sizeof(*MergerShadowAll.TmpCounter())); + runKernel({{1, -WarpSize(), 0, deviceType, RecoStep::TPCMerging}}, MergerShadow.TmpCounter(), 2 * NSECTORS * sizeof(*MergerShadow.TmpCounter())); runKernel(GetGridBlk(std::max(2u, numBlocks), 0, deviceType), 2, 3, 0); RunTPCTrackingMerger_MergeBorderTracks(GPUTPCGMMerger::mergeModes::mergeBetweenSector | GPUTPCGMMerger::mergeModes::mergeAtMidRow, deviceType); RunTPCTrackingMerger_Resolve(0, 1, deviceType); + runKernel({{1, -WarpSize(), 0, deviceType, RecoStep::TPCMerging}}, MergerShadow.TmpCounter(), 2 * NSECTORS * sizeof(*MergerShadow.TmpCounter())); // Merge between sectors - transport to the left / right edge of the sector and rotate horizontally - runKernel({{1, -WarpSize(), 0, deviceType, RecoStep::TPCMerging}}, MergerShadowAll.TmpCounter(), 2 * NSECTORS * sizeof(*MergerShadowAll.TmpCounter())); runKernel(GetGridBlk(std::max(2u, numBlocks), 0, deviceType), 0, 1, 0); RunTPCTrackingMerger_MergeBorderTracks(GPUTPCGMMerger::mergeModes::mergeBetweenSector, deviceType); RunTPCTrackingMerger_Resolve(0, 1, deviceType); + runKernel({{1, -WarpSize(), 0, deviceType, RecoStep::TPCMerging}}, MergerShadow.TmpCounter(), 2 * NSECTORS * sizeof(*MergerShadow.TmpCounter())); // Merge between sectors - use original track param - runKernel({{1, -WarpSize(), 0, deviceType, RecoStep::TPCMerging}}, MergerShadowAll.TmpCounter(), 2 * NSECTORS * sizeof(*MergerShadowAll.TmpCounter())); runKernel(GetGridBlk(std::max(2u, numBlocks), 0, deviceType), 0, 1, 1); RunTPCTrackingMerger_MergeBorderTracks(GPUTPCGMMerger::mergeModes::mergeBetweenSector | GPUTPCGMMerger::mergeModes::mergeWithOriginalParameters, deviceType); RunTPCTrackingMerger_Resolve(0, 1, deviceType); DoDebugAndDump(RecoStep::TPCMerging, GPUChainTrackingDebugFlags::TPCMergingMatching, doGPU, Merger, &GPUTPCGMMerger::DumpMergedBetweenSectors, *mDebugFile); - runKernel({{1, -WarpSize(), 0, deviceType, RecoStep::TPCMerging}}, MergerShadowAll.TmpCounter(), 2 * NSECTORS * sizeof(*MergerShadowAll.TmpCounter())); + runKernel({{1, -WarpSize(), 0, deviceType, RecoStep::TPCMerging}}, MergerShadow.TmpCounter(), 2 * NSECTORS * sizeof(*MergerShadow.TmpCounter())); runKernel(GetGridAuto(0, deviceType)); if (GetProcessingSettings().mergerSanityCheck) { @@ -177,6 +176,9 @@ int32_t GPUChainTracking::RunTPCTrackingMerger(bool synchronizeOutput) runKernel({{1, -WarpSize(), 0, deviceType}}, 1); runKernel({{1, -WarpSize(), 0, deviceType}}, 1); } + if (!doGPU && GetProcessingSettings().debug.mergerMCLabels) { + Merger.CreateMCLabels(1, 1, 0, 0); + } DoDebugAndDump(RecoStep::TPCMerging, GPUChainTrackingDebugFlags::TPCMergingCollectedTracks, doGPU, Merger, &GPUTPCGMMerger::DumpCollected, *mDebugFile); if (param().rec.tpc.mergeCE) { @@ -204,12 +206,8 @@ int32_t GPUChainTracking::RunTPCTrackingMerger(bool synchronizeOutput) Merger.CheckCollectedTracks(); } - uint32_t maxId = Merger.NMaxClusters(); - if (maxId > Merger.NMaxClusters()) { - throw std::runtime_error("mNMaxClusters too small"); - } - runKernel({{numBlocks, -ThreadCount(), 0, deviceType, RecoStep::TPCMerging}}, MergerShadowAll.SharedCount(), maxId * sizeof(*MergerShadowAll.SharedCount())); - runKernel({{numBlocks, -ThreadCount(), 0, deviceType, RecoStep::TPCMerging}}, MergerShadowAll.ClusterAttachment(), maxId * sizeof(*MergerShadowAll.ClusterAttachment())); + runKernel({{numBlocks, -ThreadCount(), 0, deviceType, RecoStep::TPCMerging}}, MergerShadow.SharedCount(), Merger.NClusters() * sizeof(*MergerShadow.SharedCount())); + runKernel({{numBlocks, -ThreadCount(), 0, deviceType, RecoStep::TPCMerging}}, MergerShadow.ClusterAttachment(), Merger.NClusters() * sizeof(*MergerShadow.ClusterAttachment())); runKernel(GetGridAuto(0, deviceType)); CondWaitEvent(waitForTransfer, &mEvents->single); runKernel(GetGridAuto(0, deviceType)); @@ -234,10 +232,30 @@ int32_t GPUChainTracking::RunTPCTrackingMerger(bool synchronizeOutput) mOutputQueue.clear(); } - runKernel(doGPU ? GetGrid(Merger.NMergedTracks(), 0) : GetGridAuto(0), mergerSortTracks ? 1 : 0); - if (param().rec.tpc.retryRefit == 1) { - runKernel(GetGridAuto(0), -1); + if (param().rec.tpc.rebuildTrackInFit) { + runKernel({{numBlocks, -ThreadCount(), 0, deviceType, RecoStep::TPCMerging}}, MergerShadow.ClusterCandidates(), Merger.NMergedTracks() * GPUTPCGeometry::NROWS * param().rec.tpc.rebuildTrackInFitClusterCandidates * sizeof(*MergerShadow.ClusterCandidates())); + } + // mRec->SetNActiveThreads(1); // For GPUTPCGMTrackParam:: GPUCA_CADEBUG: run the host kernels from here on with 1 thread, reset below + runKernel(doGPU ? GetGrid(Merger.NMergedTracks(), 0) : GetGridAuto(0), mergerSortTracks ? 1 : 0, 0); + if (param().rec.tpc.rebuildTrackInFit) { + runKernel({{numBlocks, -ThreadCount(), 0, deviceType, RecoStep::TPCMerging}}, MergerShadow.HitWeights(), Merger.NClusters() * sizeof(*MergerShadow.HitWeights())); + runKernel(doGPU ? GetGrid(Merger.NMergedTracks(), 0) : GetGridAuto(0), 0); + for (int32_t i = 0; i < param().rec.tpc.rebuildTrackInFitClusterCandidates; i++) { + runKernel(doGPU ? GetGrid(Merger.NMergedTracks(), 0) : GetGridAuto(0), i); + if (i + 1 < param().rec.tpc.rebuildTrackInFitClusterCandidates) { + runKernel(doGPU ? GetGrid(Merger.NMergedTracks(), 0) : GetGridAuto(0), i); + } + } + runKernel(doGPU ? GetGrid(Merger.NMergedTracks(), 0) : GetGridAuto(0), 0); + DoDebugAndDump(RecoStep::TPCMerging, GPUChainTrackingDebugFlags::TPCMergingRefit, Merger, &GPUTPCGMMerger::DumpInterpolatedHits, *mDebugFile); + runKernel(doGPU ? GetGrid(Merger.NMergedTracks(), 0) : GetGridAuto(0), 0); + DoDebugAndDump(RecoStep::TPCMerging, GPUChainTrackingDebugFlags::TPCMergingRefit, Merger, &GPUTPCGMMerger::DumpRebuiltTracks, *mDebugFile); + runKernel(doGPU ? GetGrid(Merger.NMergedTracks(), 0) : GetGridAuto(0), mergerSortTracks ? 1 : 0, 1); + } + if (param().rec.tpc.retryRefit) { + runKernel(GetGridAuto(0), -1, param().rec.tpc.rebuildTrackInFit); } + runKernel(GetGridAuto(0)); DoDebugAndDump(RecoStep::TPCMerging, GPUChainTrackingDebugFlags::TPCMergingRefit, Merger, &GPUTPCGMMerger::DumpRefit, *mDebugFile); @@ -246,10 +264,12 @@ int32_t GPUChainTracking::RunTPCTrackingMerger(bool synchronizeOutput) runKernel(GetGridAuto(0, deviceType)); if (param().rec.tpc.mergeLoopersAfterburner) { runKernel(doGPU ? GetGrid(Merger.NMergedTracks(), 0, deviceType) : GetGridAuto(0, deviceType)); - if (doGPU) { - TransferMemoryResourceLinkToHost(RecoStep::TPCMerging, Merger.MemoryResMemory(), 0); - SynchronizeStream(0); // TODO: could probably synchronize on an event after runKernel - } + } + if (doGPU) { + TransferMemoryResourceLinkToHost(RecoStep::TPCMerging, Merger.MemoryResMemory(), 0); + SynchronizeStream(0); // TODO: could probably synchronize on an event after runKernel + } + if (param().rec.tpc.mergeLoopersAfterburner) { runKernel(GetGridAuto(0, deviceType)); runKernel(doGPU ? GetGrid(Merger.Memory()->nLooperMatchCandidates, 0, deviceType) : GetGridAuto(0, deviceType)); DoDebugAndDump(RecoStep::TPCMerging, GPUChainTrackingDebugFlags::TPCMergingLoopers, Merger, &GPUTPCGMMerger::DumpLoopers, *mDebugFile); @@ -268,13 +288,13 @@ int32_t GPUChainTracking::RunTPCTrackingMerger(bool synchronizeOutput) throw std::runtime_error("QA Scratch buffer exceeded"); } } - GPUMemCpy(RecoStep::TPCMerging, Merger.MergedTracks(), MergerShadowAll.MergedTracks(), Merger.NMergedTracks() * sizeof(*Merger.MergedTracks()), outputStream, 0, nullptr, waitEvent); + GPUMemCpy(RecoStep::TPCMerging, Merger.MergedTracks(), MergerShadow.MergedTracks(), Merger.NMergedTracks() * sizeof(*Merger.MergedTracks()), outputStream, 0, nullptr, waitEvent); waitEvent = nullptr; if (param().dodEdxEnabled) { - GPUMemCpy(RecoStep::TPCMerging, Merger.MergedTracksdEdx(), MergerShadowAll.MergedTracksdEdx(), Merger.NMergedTracks() * sizeof(*Merger.MergedTracksdEdx()), outputStream, 0); + GPUMemCpy(RecoStep::TPCMerging, Merger.MergedTracksdEdx(), MergerShadow.MergedTracksdEdx(), Merger.NMergedTracks() * sizeof(*Merger.MergedTracksdEdx()), outputStream, 0); } - GPUMemCpy(RecoStep::TPCMerging, Merger.Clusters(), MergerShadowAll.Clusters(), Merger.NMergedTrackClusters() * sizeof(*Merger.Clusters()), outputStream, 0); - GPUMemCpy(RecoStep::TPCMerging, Merger.ClusterAttachment(), MergerShadowAll.ClusterAttachment(), Merger.NMaxClusters() * sizeof(*Merger.ClusterAttachment()), outputStream, 0); + GPUMemCpy(RecoStep::TPCMerging, Merger.Clusters(), MergerShadow.Clusters(), Merger.NMergedTrackClusters() * sizeof(*Merger.Clusters()), outputStream, 0); + GPUMemCpy(RecoStep::TPCMerging, Merger.ClusterAttachment(), MergerShadow.ClusterAttachment(), Merger.NClusters() * sizeof(*Merger.ClusterAttachment()), outputStream, 0); } if (GetProcessingSettings().outputSharedClusterMap) { TransferMemoryResourceLinkToHost(RecoStep::TPCMerging, Merger.MemoryResOutputState(), outputStream, nullptr, waitEvent); @@ -287,6 +307,8 @@ int32_t GPUChainTracking::RunTPCTrackingMerger(bool synchronizeOutput) if (GetProcessingSettings().keepDisplayMemory && !GetProcessingSettings().keepAllMemory) { TransferMemoryResourcesToHost(RecoStep::TPCMerging, &Merger, -1, true); } + DoDebugMergerNClDump(MergerShadow); + // mRec->SetNActiveThreads(-1); mRec->ReturnVolatileDeviceMemory(); mRec->PopNonPersistentMemory(RecoStep::TPCMerging, qStr2Tag("TPCMERGE")); @@ -364,7 +386,7 @@ int32_t GPUChainTracking::RunTPCTrackingMerger(bool synchronizeOutput) } if (GetProcessingSettings().debugLevel >= 2) { - GPUInfo("TPC Merger Finished (output clusters %d / input clusters %d)", Merger.NMergedTrackClusters(), Merger.NClusters()); + GPUInfo("TPC Merger Finished (output clusters %d / input clusters %d)", Merger.NMergedTrackClusters(), Merger.NSectorHits()); } return 0; } diff --git a/GPU/GPUTracking/Global/GPUErrorCodes.h b/GPU/GPUTracking/Global/GPUErrorCodes.h index a4921f478b107..b3894cf267c17 100644 --- a/GPU/GPUTracking/Global/GPUErrorCodes.h +++ b/GPU/GPUTracking/Global/GPUErrorCodes.h @@ -33,24 +33,25 @@ GPUCA_ERROR_CODE(13, ERROR_SECTORDATA_HITINROW_OVERFLOW, SectorRow, Value, Max) GPUCA_ERROR_CODE(14, ERROR_SECTORDATA_BIN_OVERFLOW, SectorRow, Value, Max) GPUCA_ERROR_CODE(15, ERROR_SECTORDATA_Z_OVERFLOW, Sector, Value) GPUCA_ERROR_CODE(16, ERROR_MERGER_HIT_OVERFLOW, Value, Max) -GPUCA_ERROR_CODE(17, ERROR_MERGER_TRACK_OVERFLOW, Value, Max) -GPUCA_ERROR_CODE(18, ERROR_COMPRESSION_ROW_HIT_OVERFLOW, SectorRow, Value, Max) -GPUCA_ERROR_CODE(19, ERROR_LOOPER_MATCH_OVERFLOW, Value, Max) -GPUCA_ERROR_CODE(20, ERROR_CF_PEAK_OVERFLOW, Sector, Value, Max) -GPUCA_ERROR_CODE(21, ERROR_CF_CLUSTER_OVERFLOW, Sector, Value, Max) -GPUCA_ERROR_CODE(22, ERROR_CF_ROW_CLUSTER_OVERFLOW, SectorRow, Value, Max) -GPUCA_ERROR_CODE(23, ERROR_CF_GLOBAL_CLUSTER_OVERFLOW, SectorRow, Value, Max) -GPUCA_ERROR_CODE(24, ERROR_DECOMPRESSION_ATTACHED_CLUSTER_OVERFLOW, SectorRow, Value, Max) -GPUCA_ERROR_CODE(25, MAX_OVERFLOW_ERROR_NUMBER) // Overflow errors are detected as errno <= MAX_OVERFLOW_ERROR_NUMBER - -GPUCA_ERROR_CODE(26, ERROR_TPCZS_INVALID_ROW, SectorRow) // Data from invalid row is skipped -GPUCA_ERROR_CODE(27, ERROR_TPCZS_INVALID_NADC, SectorCRU, SamplesInPage, SamplesWritten) // Invalid number of ADC samples in header, existing samples were decoded -GPUCA_ERROR_CODE(28, ERROR_TPCZS_INCOMPLETE_HBF, SectorCRU, PacketCount, NextPacketCount) // Part of HBF is missing, decoding incomplete -GPUCA_ERROR_CODE(29, ERROR_TPCZS_INVALID_OFFSET, SectorEndpoint, Value, Expected) // Raw page is skipped since it contains invalid payload offset -GPUCA_ERROR_CODE(30, ERROR_TPCZS_INVALID_MAGIC_WORD, Value) // ZS header contains wrong magic word -GPUCA_ERROR_CODE(31, ERROR_TPCZS_PAGE_OVERFLOW, Position, PageEnd) // Ran out of page to decode -GPUCA_ERROR_CODE(32, ERROR_TPCZS_VERSION_MISMATCH, Value, Expected) // ZS decoder received page with wrong version -GPUCA_ERROR_CODE(33, ERROR_TPCZS_UNKNOWN, ErrorCode) // Unkown or invalid error code raised in decoder -GPUCA_ERROR_CODE(33, MAX_GPUCA_ERROR_NUMBER) +GPUCA_ERROR_CODE(17, ERROR_MERGER_REBUILD_HIT_OVERFLOW, Value, Max) +GPUCA_ERROR_CODE(18, ERROR_MERGER_TRACK_OVERFLOW, Value, Max) +GPUCA_ERROR_CODE(19, ERROR_MERGER_INTERPOLATION_OVERFLOW, Value, Max) +GPUCA_ERROR_CODE(20, ERROR_COMPRESSION_ROW_HIT_OVERFLOW, SectorRow, Value, Max) +GPUCA_ERROR_CODE(21, ERROR_LOOPER_MATCH_OVERFLOW, Value, Max) +GPUCA_ERROR_CODE(22, ERROR_CF_PEAK_OVERFLOW, Sector, Value, Max) +GPUCA_ERROR_CODE(23, ERROR_CF_CLUSTER_OVERFLOW, Sector, Value, Max) +GPUCA_ERROR_CODE(24, ERROR_CF_ROW_CLUSTER_OVERFLOW, SectorRow, Value, Max) +GPUCA_ERROR_CODE(25, ERROR_CF_GLOBAL_CLUSTER_OVERFLOW, SectorRow, Value, Max) +GPUCA_ERROR_CODE(26, ERROR_DECOMPRESSION_ATTACHED_CLUSTER_OVERFLOW, SectorRow, Value, Max) +GPUCA_ERROR_CODE(26, MAX_OVERFLOW_ERROR_NUMBER) // Overflow errors are detected as errno <= MAX_OVERFLOW_ERROR_NUMBER +GPUCA_ERROR_CODE(27, ERROR_TPCZS_INVALID_ROW, SectorRow) // Data from invalid row is skipped +GPUCA_ERROR_CODE(28, ERROR_TPCZS_INVALID_NADC, SectorCRU, SamplesInPage, SamplesWritten) // Invalid number of ADC samples in header, existing samples were decoded +GPUCA_ERROR_CODE(29, ERROR_TPCZS_INCOMPLETE_HBF, SectorCRU, PacketCount, NextPacketCount) // Part of HBF is missing, decoding incomplete +GPUCA_ERROR_CODE(30, ERROR_TPCZS_INVALID_OFFSET, SectorEndpoint, Value, Expected) // Raw page is skipped since it contains invalid payload offset +GPUCA_ERROR_CODE(31, ERROR_TPCZS_INVALID_MAGIC_WORD, Value) // ZS header contains wrong magic word +GPUCA_ERROR_CODE(32, ERROR_TPCZS_PAGE_OVERFLOW, Position, PageEnd) // Ran out of page to decode +GPUCA_ERROR_CODE(33, ERROR_TPCZS_VERSION_MISMATCH, Value, Expected) // ZS decoder received page with wrong version +GPUCA_ERROR_CODE(34, ERROR_TPCZS_UNKNOWN, ErrorCode) // Unkown or invalid error code raised in decoder +GPUCA_ERROR_CODE(34, MAX_GPUCA_ERROR_NUMBER) // #define GPUCA_CHECK_TPCZS_CORRUPTION diff --git a/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx b/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx index cf8ea3c4f4675..f0183cd5472b6 100644 --- a/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx +++ b/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx @@ -17,6 +17,7 @@ // #define GPUCA_CADEBUG_ENABLED #include "GPUCommonDef.h" +#include "GPUDef.h" #if !defined(GPUCA_GPUCODE) && (defined(GPUCA_MERGER_BY_MC_LABEL) || defined(GPUCA_CADEBUG_ENABLED) || GPUCA_MERGE_LOOPER_MC) #include "AliHLTTPCClusterMCData.h" @@ -148,6 +149,7 @@ using namespace o2::gpu::internal; #include "GPUQA.h" #include "GPUMemorySizeScalers.h" +#include "GPUQAHelper.h" GPUTPCGMMerger::GPUTPCGMMerger() { @@ -164,7 +166,7 @@ GPUTPCGMMerger::GPUTPCGMMerger() } // DEBUG CODE -#if !defined(GPUCA_GPUCODE) && (defined(GPUCA_MERGER_BY_MC_LABEL) || defined(GPUCA_CADEBUG_ENABLED) || GPUCA_MERGE_LOOPER_MC) +#if defined(GPUCA_MERGER_BY_MC_LABEL) || defined(GPUCA_CADEBUG_ENABLED) || GPUCA_MERGE_LOOPER_MC #include "GPUQAHelper.h" template @@ -398,12 +400,20 @@ void* GPUTPCGMMerger::SetPointersMerger(void* mem) computePointerWithAlignment(mem, mTrackIDs, GPUTPCGeometry::NSECTORS * mNMaxSingleSectorTracks); // UnpackResetIds - RefitSectorTracks - UnpackSectorGlobal memMax = (void*)std::max((size_t)mem, (size_t)memMax); mem = memBase; - computePointerWithAlignment(mem, mTrackSort, mNMaxTracks); // PrepareForFit0 - SortTracksQPt - PrepareForFit1 - PrepareForFit1 / Finalize0 - Finalize2 - computePointerWithAlignment(mem, mSharedCount, mNMaxClusters); + computePointerWithAlignment(mem, mSharedCount, mNClusters); // PrepareForFit0 - SortTracksQPt - PrepareForFit1 - PrepareForFit1 / Finalize0 - Finalize2 + computePointerWithAlignment(mem, mTrackSort, mNMaxTracks); memMax = (void*)std::max((size_t)mem, (size_t)memMax); mem = memBase; - 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 + computePointerWithAlignment(mem, mLoopData, mNMaxTracks); // GPUTPCGMMergerTrackFit - GPUTPCGMMergerFollowLoopers, Reducing mNMaxTracks does not save memory, other parts are larger anyway + if (mRec->GetParam().rec.tpc.rebuildTrackInFit) { + computePointerWithAlignment(mem, mSharedCount2, mNClusters); + computePointerWithAlignment(mem, mClusterCandidates, mNMaxTracks * GPUTPCGeometry::NROWS * Param().rec.tpc.rebuildTrackInFitClusterCandidates); + computePointerWithAlignment(mem, mTrackRebuildHelper, mNMaxTracks); + computePointerWithAlignment(mem, mHitWeights, mNClusters); + } + if (mRec->GetParam().rec.tpc.retryRefit) { + computePointerWithAlignment(mem, mRetryRefitIds, mNMaxTracks); + } memMax = (void*)std::max((size_t)mem, (size_t)memMax); mem = memBase; computePointerWithAlignment(mem, mLooperCandidates, mNMaxLooperMatches); // MergeLoopers 1-3 @@ -430,6 +440,9 @@ void* GPUTPCGMMerger::SetPointersRefitScratch(void* mem) void* GPUTPCGMMerger::SetPointersOutput(void* mem) { computePointerWithAlignment(mem, mMergedTracks, mNMaxTracks); + if (mRec->GetProcessingSettings().debug.mergerMCLabels) { + computePointerWithAlignment(mem, mMergedTrackMC, mNMaxTracks); + } if (mRec->GetParam().dodEdxEnabled) { computePointerWithAlignment(mem, mMergedTracksdEdx, mNMaxTracks); if (mRec->GetParam().rec.tpc.dEdxClusterRejectionFlagMask != mRec->GetParam().rec.tpc.dEdxClusterRejectionFlagMaskAlt) { @@ -437,14 +450,14 @@ void* GPUTPCGMMerger::SetPointersOutput(void* mem) } } computePointerWithAlignment(mem, mClusters, mNMaxMergedTrackClusters); - computePointerWithAlignment(mem, mClusterAttachment, mNMaxClusters); + computePointerWithAlignment(mem, mClusterAttachment, mNClusters); return mem; } void* GPUTPCGMMerger::SetPointersOutputState(void* mem) { if ((mRec->GetRecoSteps() & gpudatatypes::RecoStep::Refit) || mRec->GetProcessingSettings().outputSharedClusterMap) { - computePointerWithAlignment(mem, mClusterStateExt, mNMaxClusters); + computePointerWithAlignment(mem, mClusterStateExt, mNClusters); } else { mClusterStateExt = nullptr; } @@ -497,33 +510,33 @@ void GPUTPCGMMerger::RegisterMemoryAllocation() void GPUTPCGMMerger::SetMaxData(const GPUTrackingInOutPointers& io) { mNTotalSectorTracks = 0; - mNClusters = 0; + mNSectorHits = 0; mNMaxSingleSectorTracks = 0; for (int32_t iSector = 0; iSector < NSECTORS; iSector++) { uint32_t ntrk = *mRec->GetConstantMem().tpcTrackers[iSector].NTracks(); mNTotalSectorTracks += ntrk; - mNClusters += *mRec->GetConstantMem().tpcTrackers[iSector].NTrackHits(); + mNSectorHits += *mRec->GetConstantMem().tpcTrackers[iSector].NTrackHits(); if (mNMaxSingleSectorTracks < ntrk) { mNMaxSingleSectorTracks = ntrk; } } - mNMaxMergedTrackClusters = mRec->MemoryScalers()->NTPCMergedTrackHits(mNClusters); + mNMaxMergedTrackClusters = mRec->MemoryScalers()->NTPCMergedTrackHits(mNSectorHits); if (CAMath::Abs(Param().polynomialField.GetNominalBz()) < (gpu_common_constants::kZeroFieldCut * gpu_common_constants::kCLight)) { mNMaxTracks = mRec->MemoryScalers()->getValue(mNTotalSectorTracks, mNTotalSectorTracks); // 0 magnetic field } else { mNMaxTracks = mRec->MemoryScalers()->NTPCMergedTracks(mNTotalSectorTracks); } if (io.clustersNative) { - mNMaxClusters = io.clustersNative->nClustersTotal; + mNClusters = io.clustersNative->nClustersTotal; } else if (mRec->GetRecoSteps() & gpudatatypes::RecoStep::TPCSectorTracking) { - mNMaxClusters = 0; + mNClusters = 0; for (int32_t i = 0; i < NSECTORS; i++) { - mNMaxClusters += mRec->GetConstantMem().tpcTrackers[i].NHitsTotal(); + mNClusters += mRec->GetConstantMem().tpcTrackers[i].NHitsTotal(); } } else { - mNMaxClusters = mNClusters; + mNClusters = mNSectorHits; } - mNMaxLooperMatches = mNMaxClusters / 4; // We have that much scratch memory anyway + mNMaxLooperMatches = mNClusters / 4; // We have that much scratch memory anyway } int32_t GPUTPCGMMerger::CheckSectors() @@ -539,6 +552,34 @@ int32_t GPUTPCGMMerger::CheckSectors() return 0; } +void GPUTPCGMMerger::CreateMCLabels(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread) +{ + const o2::tpc::ClusterNativeAccess* GPUrestrict() clusters = GetConstantMem()->ioPtrs.clustersNative; + if (clusters == nullptr || clusters->clustersMCTruth == nullptr) { + return; + } + if (mMergedTrackMC == nullptr) { + return; + } + + auto labelAssigner = GPUTPCTrkLbl(clusters->clustersMCTruth, 0.1f); + for (int32_t i = get_global_id(0); i < NMergedTracks(); i += get_global_size(0)) { + const auto& trk = mMergedTracks[i]; + if (!trk.OK()) { + continue; + } + labelAssigner.reset(); + for (uint32_t j = 0; j < trk.NClusters(); j++) { + const auto& cl = mClusters[trk.FirstClusterRef() + j]; + if (cl.state & GPUTPCGMMergedTrackHit::flagReject) { + continue; + } + labelAssigner.addLabel(cl.num); + } + mMergedTrackMC[i] = labelAssigner.computeLabel(); + } +} + #endif // GPUCA_GPUCODE GPUd() void GPUTPCGMMerger::ClearTrackLinks(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, bool output) @@ -566,7 +607,7 @@ GPUd() int32_t GPUTPCGMMerger::RefitSectorTrack(GPUTPCGMSectorTrack& sectorTrack trk.QPt() = inTrack->Param().GetQPt(); trk.TOffset() = Param().par.continuousTracking ? GetConstantMem()->calibObjects.fastTransform->convZOffsetToVertexTime(sector, inTrack->Param().GetZOffset(), Param().continuousMaxTimeBin) : 0; const auto tmp = sectorTrack.ClusterTN() > sectorTrack.ClusterT0() ? std::array{sectorTrack.ClusterTN(), sectorTrack.ClusterT0()} : std::array{sectorTrack.ClusterT0(), sectorTrack.ClusterTN()}; - trk.ShiftZ(this, sector, tmp[0], tmp[1], inTrack->Param().GetX()); // We do not store the inner / outer cluster X, so we just use the track X instead + trk.ShiftZ(*this, sector, tmp[0], tmp[1], inTrack->Param().GetX()); // We do not store the inner / outer cluster X, so we just use the track X instead sectorTrack.SetX2(0.f); for (int32_t way = 0; way < 2; way++) { if (way) { @@ -1657,11 +1698,11 @@ GPUd() void GPUTPCGMMerger::CollectMergedTracks(int32_t nBlocks, int32_t nThread const GPUTPCTracker& trk = GetConstantMem()->tpcTrackers[t->Sector()]; const GPUTPCHitId& ic = trk.TrackHits()[t->OrigTrack()->FirstHitID() + i]; uint32_t id = trk.Data().ClusterDataIndex(trk.Data().Row(ic.RowIndex()), ic.HitIndex()) + GetConstantMem()->ioPtrs.clustersNative->clusterOffset[t->Sector()][0]; - *c2 = trackCluster{id, (uint8_t)ic.RowIndex(), t->Sector()}; + *c2 = trackCluster{.id = id, .row = (uint8_t)ic.RowIndex(), .sector = t->Sector(), .error = 0.f, .weight = 0, .best = 0}; } nHits += nTrackHits; } - if (nHits < Param().tpcMinHitsB5(trbase->QPt() * Param().qptB5Scaler)) { + if (nHits < Param().tpcMinHitsB5Abs(CAMath::Abs(trbase->QPt() * Param().qptB5Scaler))) { break; } @@ -1838,11 +1879,13 @@ GPUd() void GPUTPCGMMerger::PrepareForFit1(int32_t nBlocks, int32_t nThreads, in GPUTPCGMMergedTrack& trk = mMergedTracks[i]; if (trk.OK()) { for (uint32_t j = 0; j < trk.NClusters(); j++) { - uint32_t weight = attachAttached | attachGood; - if (CAMath::Abs(trk.GetParam().GetQPt() * Param().qptB5Scaler) <= Param().rec.tpc.rejectQPtB5 && !trk.MergedLooper() && trk.Leg() == 0) { - weight |= attachProtect; + if (!Param().rec.tpc.rebuildTrackInFit) { + uint32_t weight = attachAttached | attachGood; + if (CAMath::Abs(trk.GetParam().GetQPt() * Param().qptB5Scaler) <= Param().rec.tpc.rejectQPtB5 && !trk.MergedLooper() && trk.Leg() == 0) { + weight |= attachProtect; + } + CAMath::AtomicMax(&mClusterAttachment[mClusters[trk.FirstClusterRef() + j].num], weight); } - CAMath::AtomicMax(&mClusterAttachment[mClusters[trk.FirstClusterRef() + j].num], weight); CAMath::AtomicAdd(&mSharedCount[mClusters[trk.FirstClusterRef() + j].num], 1u); } if (!trk.CCE() && !trk.MergedLooper()) { @@ -1851,7 +1894,7 @@ GPUd() void GPUTPCGMMerger::PrepareForFit1(int32_t nBlocks, int32_t nThreads, in const auto& GPUrestrict() cls = GetConstantMem()->ioPtrs.clustersNative->clustersLinear; float z0 = cls[cl0.num].getTime(), zn = cls[cln.num].getTime(); const auto tmp = zn > z0 ? std::array{zn, z0, GPUTPCGeometry::Row2X(cln.row)} : std::array{z0, zn, GPUTPCGeometry::Row2X(cl0.row)}; - trk.Param().ShiftZ(this, cl0.sector, tmp[0], tmp[1], tmp[2]); + trk.Param().ShiftZ(*this, cl0.sector, tmp[0], tmp[1], tmp[2]); updTrk = &trk; while (updTrk->PrevSegment() >= 0) { auto next = &mMergedTracks[updTrk->PrevSegment()]; @@ -1873,8 +1916,8 @@ GPUd() void GPUTPCGMMerger::PrepareForFit2(int32_t nBlocks, int32_t nThreads, in mClusters[i].state |= GPUTPCGMMergedTrackHit::flagShared; } } - if (mClusterStateExt) { - for (uint32_t i = iBlock * nThreads + iThread; i < mNMaxClusters; i += nBlocks * nThreads) { + if (!Param().rec.tpc.rebuildTrackInFit && mClusterStateExt) { + for (uint32_t i = iBlock * nThreads + iThread; i < mNClusters; i += nBlocks * nThreads) { uint8_t state = GetConstantMem()->ioPtrs.clustersNative->clustersLinear[i].getFlags(); if (mSharedCount[i] > 1) { state |= GPUTPCGMMergedTrackHit::flagShared; @@ -1925,7 +1968,7 @@ GPUd() void GPUTPCGMMerger::Finalize1(int32_t nBlocks, int32_t nThreads, int32_t GPUd() void GPUTPCGMMerger::Finalize2(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread) { - for (uint32_t i = iBlock * nThreads + iThread; i < mNMaxClusters; i += nThreads * nBlocks) { + for (uint32_t i = iBlock * nThreads + iThread; i < mNClusters; i += nThreads * nBlocks) { if (mClusterAttachment[i] != 0) { mClusterAttachment[i] = (mClusterAttachment[i] & attachFlagMask) | mTrackSort[mClusterAttachment[i] & attachTrackMask]; } @@ -1965,7 +2008,7 @@ GPUd() void GPUTPCGMMerger::MergeLoopersInit(int32_t nBlocks, int32_t nThreads, for (uint32_t k = 0;k < trk.NClusters();k++) { float xx, yy, zz; const ClusterNative& GPUrestrict() cl = GetConstantMem()->ioPtrs.clustersNative->clustersLinear[mClusters[trk.FirstClusterRef() + k].num]; - GetConstantMem()->calibObjects.fastTransformHelper->Transform(mClusters[trk.FirstClusterRef() + k].sector, mClusters[trk.FirstClusterRef() + k].row, cl.getPad(), cl.getTime(), xx, yy, zz, p.GetTOffset()); + GetConstantMem()->calibObjects.fastTransform->Transform(mClusters[trk.FirstClusterRef() + k].sector, mClusters[trk.FirstClusterRef() + k].row, cl.getPad(), cl.getTime(), xx, yy, zz, p.GetTOffset()); float sa2, ca2; CAMath::SinCos(Param().Alpha(mClusters[trk.FirstClusterRef() + k].sector), sa2, ca2); float cx = ca2 * xx - sa2 * yy; @@ -2120,3 +2163,180 @@ GPUd() void GPUTPCGMMerger::MergeLoopersMain(int32_t nBlocks, int32_t nThreads, } #endif*/ } + +GPUd() uint32_t GPUTPCGMMerger::MinHitsForQPt(float qPt, float err2QPt) const +{ + return Param().tpcMinHitsB5Abs((CAMath::Abs(qPt) + CAMath::Min(CAMath::Sqrt(CAMath::Max(0.f, err2QPt)), Param().rec.tpc.trackFitMinHitsMaxQPtIncrease * CAMath::Abs(qPt))) * Param().qptB5Scaler); +} + +GPUd() void GPUTPCGMMerger::PrepareHitWeights(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread) +{ + for (uint32_t i = iBlock * nThreads + iThread; i < mNClusters; i += nBlocks * nThreads) { + mSharedCount2[i] = 0; + } + if (iThread == 0 && iBlock == 0) { + mMemory->nMergedTrackClusters = 0; + } + for (int32_t i = get_global_id(0); i < NMergedTracks(); i += get_global_size(0)) { + const auto& trk = mMergedTracks[i]; + if (!trk.OK()) { + continue; + } + mTrackRebuildHelper[i].reverse = mClusters[trk.FirstClusterRef()].row < mClusters[trk.FirstClusterRef() + trk.NClusters() - 1].row; + if (!(Param().rec.tpc.disableRebuildAttachment & 8)) { + const uint32_t nMaxCandidates = Param().rec.tpc.rebuildTrackInFitClusterCandidates; + int lastRow = -1; + for (uint32_t j = 0; j < trk.NClusters(); j++) { + const auto& cl = mClusters[trk.FirstClusterRef() + j]; + auto* candidates = &mClusterCandidates[(i * GPUTPCGeometry::NROWS + cl.row) * nMaxCandidates]; + if ((!(cl.state & GPUTPCGMMergedTrackHit::flagReject) || trk.GetParam().GetNDF() <= 0) && (candidates[0].id == 0 || cl.row == lastRow)) { // Row not searched: the original clusters of the row compete as normal candidates with fixed error + for (uint32_t k = 0; k < nMaxCandidates; k++) { + if (candidates[k].id == 0) { + candidates[k] = {.id = cl.num + 2, .row = cl.row, .sector = cl.sector, .error = Param().rec.tpc.rebuildTrackFallbackErr, .weight = 0, .best = 0}; + break; + } + } + lastRow = cl.row; + } + } + } + } +} + +GPUd() void GPUTPCGMMerger::ComputeHitWeights(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, int32_t iteration) +{ + for (int32_t i = get_global_id(0); i < NMergedTracks(); i += get_global_size(0)) { + const auto& trk = mMergedTracks[i]; + if (!trk.OK()) { + continue; + } + for (uint32_t j = 0; j < GPUTPCGeometry::NROWS; j++) { + auto* candidates = &mClusterCandidates[(i * GPUTPCGeometry::NROWS + j) * Param().rec.tpc.rebuildTrackInFitClusterCandidates]; + if (candidates[0].best == 0) { + if (candidates[iteration].id >= 2 && (iteration == 0 || candidates[iteration].error <= CAMath::Max(Param().rec.tpc.rebuildTrackAltCandidateGoodErr, candidates[0].error * Param().rec.tpc.rebuildTrackAltCandidateMaxErrFactor))) { // Don't fall back to a much worse candidate, the best one is then rather a shared cluster + float weightf = (trk.GetParam().GetNDF() > 0 ? (trk.GetParam().GetNDF() + 5) * trk.GetParam().GetNDF() * 64 / trk.GetParam().GetChi2() : (2 * (int32_t)trk.NClusters() - 5) * 128 / Param().rec.tpc.rebuildTrackFallbackTrackChi2) / (1.f + Param().rec.tpc.rebuildErrorWeightFactor * candidates[iteration].error); // Tracks with NDF <= 0 use their nominal number of clusters and assume a fixed chi2 per cluster + int32_t weight = weightf > 0.f ? (weightf > 2e9f ? 2e9f : (int32_t)weightf) : 0; + candidates[iteration].weight = weight; + CAMath::AtomicMax(&mHitWeights[candidates[iteration].id - 2], weight); + CADEBUG(printf("REBUILD: iTrk %d Iteration %d Row %d Cluster %d Weight %d\n", i, iteration, j, candidates[iteration].id - 2, weight)); + } else { + candidates[0].best = -2; + } + } + } + } +} + +GPUd() void GPUTPCGMMerger::ResolveHitWeights1(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, int32_t iteration) +{ + for (int32_t i = get_global_id(0); i < NMergedTracks(); i += get_global_size(0)) { + auto& trk = mMergedTracks[i]; + if (!trk.OK()) { + continue; + } + for (unsigned int j = 0; j < GPUTPCGeometry::NROWS; j++) { + auto* candidates = &mClusterCandidates[(i * GPUTPCGeometry::NROWS + j) * Param().rec.tpc.rebuildTrackInFitClusterCandidates]; + if (candidates[0].best == 0) { + if (candidates[iteration].weight == mHitWeights[candidates[iteration].id - 2]) { + candidates[0].best = iteration + 1; + } + } + } + } +} + +GPUd() void GPUTPCGMMerger::ResolveHitWeights2(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread) +{ + int32_t nIter = Param().rec.tpc.rebuildTrackInFitClusterCandidates; + for (int32_t i = get_global_id(0); i < NMergedTracks(); i += get_global_size(0)) { + auto& trk = mMergedTracks[i]; + if (!trk.OK()) { + trk.SetNClusters(0); + continue; + } + uint32_t nHits = 0; + uint8_t sharedRowCandidates[GPUTPCGeometry::NROWS]; + uint8_t nSharedRowCandidates = 0; + for (unsigned int j = 0; j < GPUTPCGeometry::NROWS; j++) { + auto* candidates = &mClusterCandidates[(i * GPUTPCGeometry::NROWS + j) * Param().rec.tpc.rebuildTrackInFitClusterCandidates]; + if (candidates[0].best == 0 && candidates[nIter - 1].weight == mHitWeights[candidates[nIter - 1].id - 2]) { + candidates[0].best = nIter; + } + if (candidates[0].best > 0) { + nHits++; + if (candidates[candidates[0].best - 1].weight != mHitWeights[candidates[candidates[0].best - 1].id - 2]) { + mSharedCount2[candidates[candidates[0].best - 1].id - 2] = 1; // shared between tracks of different iterations + } + } else if (candidates[0].best != -1 && candidates[0].id >= 2) { + sharedRowCandidates[nSharedRowCandidates++] = j; + } + } + + if (nSharedRowCandidates) { + const uint32_t nShared = CAMath::Min(nSharedRowCandidates, CAMath::Float2UIntRn(nHits * Param().rec.tpc.rebuildTrackMaxSharedFraction)); + if (nShared) { + const uint32_t n = Param().rec.tpc.rebuildTrackInFitClusterCandidates; + const auto* candidates = &mClusterCandidates[i * GPUTPCGeometry::NROWS * n]; + GPUCommonAlgorithm::sort(sharedRowCandidates, sharedRowCandidates + nSharedRowCandidates, [n, candidates](const uint8_t& a, const uint8_t& b) { + const auto& ca = candidates[a * n]; + const auto& cb = candidates[b * n]; + return GPUCA_DETERMINISTIC_CODE(ca.error == cb.error ? ca.id < cb.id : ca.error < cb.error, ca.error < cb.error); + }); + for (uint32_t j = 0; j < nShared; j++) { + auto* cj = &mClusterCandidates[(i * GPUTPCGeometry::NROWS + sharedRowCandidates[j]) * Param().rec.tpc.rebuildTrackInFitClusterCandidates]; + CADEBUG(printf("REBUILD: iTrk %d Recovering shared %d / %d: row %d %d, error %f\n", i, j, nShared, (int)sharedRowCandidates[j], cj[0].id - 2, cj[0].error)); + cj[0].best = 1; + mSharedCount2[cj[0].id - 2] = 1; // TODO: Should we also recover iter0 clusters? + } + nHits += nShared; + } + } + + const uint32_t iOut = CAMath::AtomicAdd(&mMemory->nMergedTrackClusters, nHits); + if (iOut + nHits > mNMaxMergedTrackClusters) { + raiseError(GPUErrors::ERROR_MERGER_REBUILD_HIT_OVERFLOW, iOut, mNMaxMergedTrackClusters); + CAMath::AtomicExch(&mMemory->nMergedTrackClusters, mNMaxMergedTrackClusters); + trk.SetNClusters(0); + trk.SetOK(false); + continue; + } + + trk.SetNClusters(nHits); + trk.SetFirstClusterRef(iOut); + uint32_t written = 0; + const bool reverse = mTrackRebuildHelper[i].reverse; // TODO: CE tracks may have clusters on both sides in the rows around the crossing: make sure the track crosses the CE only once, by cutting at the best row and removing the clusters on the wrong side of the cut + for (unsigned int j = 0; j < GPUTPCGeometry::NROWS; j++) { + const auto& candidates = &mClusterCandidates[(i * GPUTPCGeometry::NROWS + j) * Param().rec.tpc.rebuildTrackInFitClusterCandidates]; + if (candidates[0].best > 0) { + auto& outCl = mClusters[iOut + (reverse ? (written) : (nHits - 1 - written))]; + const auto& best = candidates[candidates[0].best - 1]; + const ClusterNative& GPUrestrict() cl = GetConstantMem()->ioPtrs.clustersNative->clustersLinear[best.id - 2]; + uint8_t flags = (uint8_t)(cl.getFlags() & GPUTPCGMMergedTrackHit::clustererAndSharedFlags); + if ((mTrackRebuildHelper[i].highInclRowLow != 255 && j <= mTrackRebuildHelper[i].highInclRowLow) || (mTrackRebuildHelper[i].highInclRowHigh != 255 && j >= mTrackRebuildHelper[i].highInclRowHigh)) { + flags |= GPUTPCGMMergedTrackHit::flagHighIncl; + } + outCl = {.num = best.id - 2, .sector = best.sector, .row = (uint8_t)j, .state = flags}; + written++; + CADEBUG(printf("REBUILD: iTrk %d Assigned Cluster Row %d Hit %d\n", i, j, best.id - 2)); + } + } + } +} + +GPUd() void GPUTPCGMMerger::ResolveHitWeightsShared(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread) +{ + for (uint32_t i = iBlock * nThreads + iThread; i < mMemory->nMergedTrackClusters; i += nBlocks * nThreads) { + if (mSharedCount2[mClusters[i].num]) { + mClusters[i].state |= GPUTPCGMMergedTrackHit::flagShared; // TODO: What is the best criterion to set the shared flag? Perhaps already when iter0 best hit would be shared? + } + } + if (mClusterStateExt) { + for (uint32_t i = iBlock * nThreads + iThread; i < mNClusters; i += nBlocks * nThreads) { + uint8_t state = GetConstantMem()->ioPtrs.clustersNative->clustersLinear[i].getFlags(); + if (mSharedCount2[i]) { + state |= GPUTPCGMMergedTrackHit::flagShared; + } + mClusterStateExt[i] = state; + } + } +} diff --git a/GPU/GPUTracking/Merger/GPUTPCGMMerger.h b/GPU/GPUTracking/Merger/GPUTPCGMMerger.h index eefaf6fb74497..1ca312165f17a 100644 --- a/GPU/GPUTracking/Merger/GPUTPCGMMerger.h +++ b/GPU/GPUTracking/Merger/GPUTPCGMMerger.h @@ -79,10 +79,19 @@ class GPUTPCGMMerger : public GPUProcessor GPUAtomic(uint32_t) nLooperMatchCandidates; }; - struct trackCluster { + struct trackCluster { // TODO: Reduce size of this struct! uint32_t id; uint8_t row; uint8_t sector; + float error; + int32_t weight; + int32_t best; + }; + + struct trackRebuildHelper { + uint8_t reverse; + uint8_t highInclRowLow; + uint8_t highInclRowHigh; }; struct tmpSort { @@ -114,18 +123,22 @@ class GPUTPCGMMerger : public GPUProcessor GPUhdi() int32_t NMergedTracks() const { return mMemory->nMergedTracks; } GPUhdi() const GPUTPCGMMergedTrack* MergedTracks() const { return mMergedTracks; } + GPUhdi() const o2::MCCompLabel* MergedTrackMC() const { return mMergedTrackMC; } GPUhdi() GPUTPCGMMergedTrack* MergedTracks() { return mMergedTracks; } GPUhdi() const GPUdEdxInfo* MergedTracksdEdx() const { return mMergedTracksdEdx; } GPUhdi() GPUdEdxInfo* MergedTracksdEdx() { return mMergedTracksdEdx; } GPUhdi() const GPUdEdxInfo* MergedTracksdEdxAlt() const { return mMergedTracksdEdxAlt; } GPUhdi() GPUdEdxInfo* MergedTracksdEdxAlt() { return mMergedTracksdEdxAlt; } + GPUhdi() uint32_t NSectorHits() const { return mNSectorHits; } GPUhdi() uint32_t NClusters() const { return mNClusters; } - GPUhdi() uint32_t NMaxClusters() const { return mNMaxClusters; } GPUhdi() uint32_t NMaxTracks() const { return mNMaxTracks; } GPUhdi() uint32_t NMaxMergedTrackClusters() const { return mNMaxMergedTrackClusters; } GPUhdi() uint32_t NMergedTrackClusters() const { return mMemory->nMergedTrackClusters; } GPUhdi() const GPUTPCGMMergedTrackHit* Clusters() const { return mClusters; } - GPUhdi() GPUTPCGMMergedTrackHit* Clusters() { return (mClusters); } + GPUhdi() GPUTPCGMMergedTrackHit* Clusters() { return mClusters; } + GPUhdi() trackCluster* ClusterCandidates() { return mClusterCandidates; } + GPUhdi() trackRebuildHelper* TrackRebuildHelper() { return mTrackRebuildHelper; } + GPUhdi() int32_t* HitWeights() { return mHitWeights; } GPUhdi() GPUAtomic(uint32_t) * ClusterAttachment() const { return mClusterAttachment; } GPUhdi() uint32_t* TrackOrderAttach() const { return mTrackOrderAttach; } GPUhdi() uint32_t* TrackOrderProcess() const { return mTrackOrderProcess; } @@ -139,6 +152,7 @@ class GPUTPCGMMerger : public GPUProcessor GPUhdi() tmpSort* TrackSortO2() { return mTrackSortO2; } GPUhdi() internal::MergeLooperParam* LooperCandidates() { return mLooperCandidates; } GPUhdi() GPUAtomic(uint32_t) * SharedCount() { return mSharedCount; } + GPUhdi() uint8_t* SharedCount2() { return mSharedCount2; } GPUhdi() gputpcgmmergertypes::GPUTPCGMBorderRange* BorderRange(int32_t i) { return mBorderRange[i]; } GPUhdi() const gputpcgmmergertypes::GPUTPCGMBorderRange* BorderRange(int32_t i) const { return mBorderRange[i]; } GPUhdi() GPUTPCGMBorderTrack* BorderTracks(int32_t i) { return mBorder[i]; } @@ -196,6 +210,13 @@ class GPUTPCGMMerger : public GPUProcessor GPUd() void MergeLoopersInit(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread); GPUd() void MergeLoopersSort(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread); GPUd() void MergeLoopersMain(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread); + GPUd() void PrepareHitWeights(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread); + GPUd() void ComputeHitWeights(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, int32_t iteration); + GPUd() void ResolveHitWeights1(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, int32_t iteration); + GPUd() void ResolveHitWeights2(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread); + GPUd() void ResolveHitWeightsShared(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread); + GPUd() uint32_t MinHitsForQPt(float qPt, float err2QPt) const; // QPt-dependent minimum number of hits, for QPt increased by 1 sigma (at most by the fraction trackFitMinHitsMaxQPtIncrease) + GPUd() void CreateMCLabels(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread); #ifndef GPUCA_GPUCODE void DumpSectorTracks(std::ostream& out) const; @@ -209,6 +230,8 @@ class GPUTPCGMMerger : public GPUProcessor void DumpRefit(std::ostream& out) const; void DumpFinal(std::ostream& out) const; void DumpLoopers(std::ostream& out) const; + void DumpInterpolatedHits(std::ostream& out) const; + void DumpRebuiltTracks(std::ostream& out) const; void DumpTrackParam(std::ostream& out) const; void DumpTrackClusters(std::ostream& out, bool non0StateOnly = false, bool noNDF0 = false) const; @@ -261,7 +284,7 @@ class GPUTPCGMMerger : public GPUProcessor uint32_t mNMaxTracks = 0; // maximum number of output tracks uint32_t mNMaxSingleSectorTracks = 0; // max N tracks in one sector uint32_t mNMaxMergedTrackClusters = 0; // max number of clusters in output tracks (double-counting shared clusters) - uint32_t mNMaxClusters = 0; // max total unique clusters (in event) + uint32_t mNClusters = 0; // max total unique clusters (in event) uint32_t mNMaxLooperMatches = 0; // Maximum number of candidate pairs for looper matching uint16_t mMemoryResMemory = (uint16_t)-1; @@ -272,8 +295,13 @@ class GPUTPCGMMerger : public GPUProcessor uint16_t mMemoryResOutputO2MC = (uint16_t)-1; uint16_t mMemoryResOutputO2Scratch = (uint16_t)-1; - int32_t mNClusters = 0; // Total number of incoming clusters (from sector tracks) + int32_t mNSectorHits = 0; // Total number of incoming clusters (from sector tracks) GPUTPCGMMergedTrack* mMergedTracks = nullptr; //* array of output merged tracks + o2::MCCompLabel* mMergedTrackMC = nullptr; + trackCluster* mClusterCandidates = nullptr; + trackRebuildHelper* mTrackRebuildHelper = nullptr; + int32_t* mHitWeights = nullptr; + GPUdEdxInfo* mMergedTracksdEdx = nullptr; //* dEdx information GPUdEdxInfo* mMergedTracksdEdxAlt = nullptr; //* dEdx alternative information GPUTPCGMSectorTrack* mSectorTrackInfos = nullptr; //* additional information for sector tracks @@ -294,6 +322,7 @@ class GPUTPCGMMerger : public GPUProcessor uint32_t* mTrackSort = nullptr; tmpSort* mTrackSortO2 = nullptr; GPUAtomic(uint32_t) * mSharedCount = nullptr; // Must be uint32_t unfortunately for atomic support + uint8_t* mSharedCount2 = nullptr; GPUTPCGMBorderTrack* mBorderMemory = nullptr; // memory for border tracks GPUTPCGMBorderTrack* mBorder[2 * NSECTORS]; gputpcgmmergertypes::GPUTPCGMBorderRange* mBorderRangeMemory = nullptr; // memory for border tracks diff --git a/GPU/GPUTracking/Merger/GPUTPCGMMergerDump.cxx b/GPU/GPUTracking/Merger/GPUTPCGMMergerDump.cxx index 9f6858df4d3f2..d22fd2e05e0a0 100644 --- a/GPU/GPUTracking/Merger/GPUTPCGMMergerDump.cxx +++ b/GPU/GPUTracking/Merger/GPUTPCGMMergerDump.cxx @@ -150,6 +150,15 @@ void GPUTPCGMMerger::DumpTrackParam(std::ostream& out) const out << std::setprecision(ss); } +void GPUTPCGMMerger::DumpRebuiltTracks(std::ostream& out) const +{ + out << "\nTPC Merger Rebuilt Tracks\n"; + out << " Cluster Attachment\n"; + DumpTrackClusters(out, false, true); + out << " Track Params\n"; + DumpTrackParam(out); +} + void GPUTPCGMMerger::DumpMergeCE(std::ostream& out) const { DumpTrackLinks(out, true, " for CE merging"); @@ -192,7 +201,7 @@ void GPUTPCGMMerger::DumpFitPrepare(std::ostream& out) const out << " Track Clusters"; DumpTrackClusters(out); uint32_t j = 0; - for (uint32_t i = 0; i < mNMaxClusters; i++) { + for (uint32_t i = 0; i < mNClusters; i++) { if ((mClusterAttachment[i] & attachFlagMask) != 0) { if (j++ % 10 == 0) { out << "\n Cluster attachment "; @@ -241,7 +250,7 @@ void GPUTPCGMMerger::DumpFinal(std::ostream& out) const out << "Track Clusters\n"; DumpTrackClusters(out, true); uint32_t j = 0; - for (uint32_t i = 0; i < mNMaxClusters; i++) { + for (uint32_t i = 0; i < mNClusters; i++) { if ((mClusterAttachment[i] & attachFlagMask) != 0) { if (++j % 10 == 0) { out << " Cluster attachment "; @@ -255,6 +264,34 @@ void GPUTPCGMMerger::DumpFinal(std::ostream& out) const out << "\n"; } +void GPUTPCGMMerger::DumpInterpolatedHits(std::ostream& out) const +{ + out << "\nTPC Merger Interpolated Hits\n"; + for (uint32_t i = 0; i < mMemory->nMergedTracks; i++) { + const auto& trk = mMergedTracks[i]; + if (trk.OK() && trk.GetParam().GetNDF() >= 0 && trk.NClusters()) { + out << "Track " << i << ":"; + for (uint32_t j = 0; j < GPUTPCGeometry::NROWS; j++) { + auto* candidates = &mClusterCandidates[(i * GPUTPCGeometry::NROWS + j) * Param().rec.tpc.rebuildTrackInFitClusterCandidates]; + if (candidates[0].id) { + out << " Row " << j << ": "; + if (candidates[0].best) { + out << " Best " << candidates[0].best << " - "; + } + for (uint32_t k = 0; k < Param().rec.tpc.rebuildTrackInFitClusterCandidates; k++) { + if (candidates[k].id) { + out << k << ": id " << candidates[k].id << " err " << candidates[k].error << " weight " << candidates[k].weight << " - "; + } + } + out << "; "; + } + } + out << "\n"; + } + } + out << "\n"; +} + template inline void GPUTPCGMMerger::MergedTrackStreamerInternal(const GPUTPCGMBorderTrack& b1, const GPUTPCGMBorderTrack& b2, const char* name, int32_t sector1, int32_t sector2, uint8_t mergeMode, float weight, float frac) const { diff --git a/GPU/GPUTracking/Merger/GPUTPCGMMergerGPU.cxx b/GPU/GPUTracking/Merger/GPUTPCGMMergerGPU.cxx index bf01ea234a645..2ef0d0ba59748 100644 --- a/GPU/GPUTracking/Merger/GPUTPCGMMergerGPU.cxx +++ b/GPU/GPUTracking/Merger/GPUTPCGMMergerGPU.cxx @@ -19,12 +19,12 @@ using namespace o2::gpu; template <> -GPUdii() void GPUTPCGMMergerTrackFit::Thread<0>(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& smem, processorType& GPUrestrict() merger, int32_t mode) +GPUdii() void GPUTPCGMMergerTrackFit::Thread<0>(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& smem, processorType& GPUrestrict() merger, int32_t mode, int32_t rebuilt) { const int32_t iEnd = mode == -1 ? merger.Memory()->nRetryRefit : merger.NMergedTracks(); GPUCA_TBB_KERNEL_LOOP(merger.GetRec(), int32_t, ii, iEnd, { const int32_t i = mode == -1 ? merger.RetryRefitIds()[ii] : (mode ? merger.TrackOrderProcess()[ii] : ii); - GPUTPCGMTrackParam::RefitTrack(merger.MergedTracks()[i], i, &merger, mode == -1); + GPUTPCGMTrackParam::RefitTrack(merger.MergedTracks()[i], i, merger, rebuilt, mode == -1); }); } @@ -32,7 +32,7 @@ template <> GPUdii() void GPUTPCGMMergerFollowLoopers::Thread<0>(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& smem, processorType& GPUrestrict() merger) { GPUCA_TBB_KERNEL_LOOP(merger.GetRec(), uint32_t, i, merger.Memory()->nLoopData, { - GPUTPCGMTrackParam::PropagateLooper(&merger, i); + GPUTPCGMTrackParam::PropagateLooper(merger, i); }); } @@ -215,3 +215,33 @@ GPUdii() void GPUTPCGMMergerMergeLoopers::Thread<2>(int32_t nBlocks, int32_t nTh { merger.MergeLoopersMain(nBlocks, nThreads, iBlock, iThread); } + +template <> +GPUdii() void GPUTPCGMMergerHitWeights::Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& smem, processorType& GPUrestrict() merger, int32_t iteration) +{ + merger.PrepareHitWeights(nBlocks, nThreads, iBlock, iThread); +} + +template <> +GPUdii() void GPUTPCGMMergerHitWeights::Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& smem, processorType& GPUrestrict() merger, int32_t iteration) +{ + merger.ComputeHitWeights(nBlocks, nThreads, iBlock, iThread, iteration); +} + +template <> +GPUdii() void GPUTPCGMMergerHitWeights::Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& smem, processorType& GPUrestrict() merger, int32_t iteration) +{ + merger.ResolveHitWeights1(nBlocks, nThreads, iBlock, iThread, iteration); +} + +template <> +GPUdii() void GPUTPCGMMergerHitWeights::Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& smem, processorType& GPUrestrict() merger, int32_t iteration) +{ + merger.ResolveHitWeights2(nBlocks, nThreads, iBlock, iThread); +} + +template <> +GPUdii() void GPUTPCGMMergerHitWeights::Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& smem, processorType& GPUrestrict() merger, int32_t iteration) +{ + merger.ResolveHitWeightsShared(nBlocks, nThreads, iBlock, iThread); +} diff --git a/GPU/GPUTracking/Merger/GPUTPCGMMergerGPU.h b/GPU/GPUTracking/Merger/GPUTPCGMMergerGPU.h index d703f1b890443..82605f47a5f27 100644 --- a/GPU/GPUTracking/Merger/GPUTPCGMMergerGPU.h +++ b/GPU/GPUTracking/Merger/GPUTPCGMMergerGPU.h @@ -36,7 +36,7 @@ class GPUTPCGMMergerTrackFit : public GPUTPCGMMergerGeneral { public: template - GPUd() static void Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& smem, processorType& merger, int32_t mode); + GPUd() static void Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& smem, processorType& merger, int32_t mode, int32_t rebuilt); }; class GPUTPCGMMergerFollowLoopers : public GPUTPCGMMergerGeneral @@ -180,6 +180,19 @@ class GPUTPCGMMergerMergeLoopers : public GPUTPCGMMergerGeneral GPUd() static void Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& smem, processorType& merger); }; +class GPUTPCGMMergerHitWeights : public GPUTPCGMMergerGeneral +{ + public: + enum K { defaultKernel = 0, + prepare = 0, + compute = 1, + resolve1 = 2, + resolve2 = 3, + resolveShared = 4 }; + template + GPUd() static void Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& smem, processorType& merger, int32_t iteration); +}; + } // namespace o2::gpu #endif diff --git a/GPU/GPUTracking/Merger/GPUTPCGMMergerTypes.h b/GPU/GPUTracking/Merger/GPUTPCGMMergerTypes.h index 99c47b6228889..e35c697c6ac49 100644 --- a/GPU/GPUTracking/Merger/GPUTPCGMMergerTypes.h +++ b/GPU/GPUTracking/Merger/GPUTPCGMMergerTypes.h @@ -34,10 +34,13 @@ enum attachTypes { attachProtect = 0x80000000, struct InterpolationErrorHit { float posY, posZ; GPUCA_PAR_MERGER_INTERPOLATION_ERROR_TYPE_A errorY, errorZ; + GPUdi() bool isValid() const { return errorY >= (GPUCA_PAR_MERGER_INTERPOLATION_ERROR_TYPE_A)0; } + GPUdi() void markInvalid() { errorY = (GPUCA_PAR_MERGER_INTERPOLATION_ERROR_TYPE_A)-1; } }; struct InterpolationErrors { - InterpolationErrorHit hit[constants::MERGER_MAX_TRACK_CLUSTERS]; + static constexpr size_t size = constants::MERGER_MAX_TRACK_CLUSTERS; + InterpolationErrorHit hit[size]; }; struct GPUResolveSharedMemory : public GPUKernelTemplate::GPUSharedMemoryScan64 { diff --git a/GPU/GPUTracking/Merger/GPUTPCGMO2Output.cxx b/GPU/GPUTracking/Merger/GPUTPCGMO2Output.cxx index 4bcffd8a31c46..b2009aba09d41 100644 --- a/GPU/GPUTracking/Merger/GPUTPCGMO2Output.cxx +++ b/GPU/GPUTracking/Merger/GPUTPCGMO2Output.cxx @@ -76,7 +76,7 @@ GPUdii() void GPUTPCGMO2Output::Thread(int32_t nBlock if (nCl == 0) { continue; } - if (nCl + 2 < merger.Param().tpcMinHitsB5(tracks[i].GetParam().GetQPt() * merger.Param().qptB5Scaler)) { // Give 2 hits tolerance in the primary leg, compared to the full fit of the looper + if (nCl + 2 < merger.MinHitsForQPt(tracks[i].GetParam().GetQPt(), tracks[i].GetParam().GetErr2QPt())) { // Give 2 hits tolerance in the primary leg, compared to the full fit of the looper continue; } if (merger.Param().rec.tpc.minNClustersFinalTrack != -1 && nCl < (uint32_t)merger.Param().rec.tpc.minNClustersFinalTrack) { @@ -215,9 +215,11 @@ GPUdii() void GPUTPCGMO2Output::Thread(int32_t nBlocks if (track.PrevSegment() >= 0) { const GPUTPCGMMergedTrack* chkTrk = track.GetFirstSegment(tracks, merger.Param().rec.enableCyclicGraphWorkarounds); - const auto& firstPrevCluster = trackClusters[chkTrk->FirstClusterRef()]; - t1 = clusters->clustersLinear[firstPrevCluster.num].getTime(); - sector1 = firstPrevCluster.sector; + if (chkTrk->NClusters()) { // First segment can have lost all its clusters in the rebuild, then its cluster reference is invalid + const auto& firstPrevCluster = trackClusters[chkTrk->FirstClusterRef()]; + t1 = clusters->clustersLinear[firstPrevCluster.num].getTime(); + sector1 = firstPrevCluster.sector; + } } bool cce = track.CCE() && ((sector1 < MAXSECTOR / 2) ^ (sector2 < MAXSECTOR / 2)); diff --git a/GPU/GPUTracking/Merger/GPUTPCGMPropagator.cxx b/GPU/GPUTracking/Merger/GPUTPCGMPropagator.cxx index 974a0b3ea2074..ddfefe8636af5 100644 --- a/GPU/GPUTracking/Merger/GPUTPCGMPropagator.cxx +++ b/GPU/GPUTracking/Merger/GPUTPCGMPropagator.cxx @@ -608,7 +608,7 @@ GPUd() float GPUTPCGMPropagator::PredictChi2(float posY, float posZ, float err2Y } } -GPUd() int32_t GPUTPCGMPropagator::Update(float posY, float posZ, int32_t iRow, const GPUParam& GPUrestrict() param, int16_t clusterState, int8_t rejectChi2, bool refit, int8_t sector, float time, float avgInvCharge, float invCharge) +GPUd() int32_t GPUTPCGMPropagator::Update(float posY, float posZ, int32_t iRow, const GPUParam& GPUrestrict() param, int16_t clusterState, bool rejectChi2, bool refit, int8_t sector, float time, float avgInvCharge, float invCharge) { float err2Y, err2Z; GetErr2(err2Y, err2Z, param, posZ, iRow, clusterState, sector, time, avgInvCharge, invCharge); @@ -616,7 +616,7 @@ GPUd() int32_t GPUTPCGMPropagator::Update(float posY, float posZ, int32_t iRow, return Update(posY, posZ, iRow, param, clusterState, rejectChi2, refit, err2Y, err2Z); } -GPUd() int32_t GPUTPCGMPropagator::Update(float posY, float posZ, int32_t iRow, const GPUParam& GPUrestrict() param, int16_t clusterState, int8_t rejectChi2, bool refit, float err2Y, float err2Z) +GPUd() int32_t GPUTPCGMPropagator::Update(float posY, float posZ, int32_t iRow, const GPUParam& GPUrestrict() param, int16_t clusterState, bool rejectChi2, bool refit, float err2Y, float err2Z) { if (mT->NDF() == -5) { // first measurement: no need to filter, as the result is known in advance. just set it. mT->ResetCovariance(); @@ -635,80 +635,83 @@ GPUd() int32_t GPUTPCGMPropagator::Update(float posY, float posZ, int32_t iRow, return 0; } - return Update(posY, posZ, clusterState, rejectChi2 == rejectDirect, err2Y, err2Z, ¶m); + return Update(posY, posZ, clusterState, rejectChi2, err2Y, err2Z, ¶m); } -GPUd() int32_t GPUTPCGMPropagator::InterpolateReject(const GPUParam& GPUrestrict() param, float posY, float posZ, int16_t clusterState, int8_t rejectChi2, gputpcgmmergertypes::InterpolationErrorHit* inter, float err2Y, float err2Z, float deltaZ) +GPUd() void GPUTPCGMPropagator::InterpolateFill(gputpcgmmergertypes::InterpolationErrorHit* inter) { float* GPUrestrict() mC = mT->Cov(); float* GPUrestrict() mP = mT->Par(); - if (rejectChi2 == rejectInterFill) { - inter->posY = mP[0]; - inter->posZ = mP[1]; - if (mT->NDF() <= 0) { - inter->errorY = inter->errorZ = 100.f; - } else { - inter->errorY = mC[0]; - inter->errorZ = mC[2]; - } - } else if (rejectChi2 == rejectInterReject) { - float chi2Y, chi2Z; - if (mT->NDF() <= 0) { - chi2Y = CAMath::Square((float)inter->posY - posY) / ((float)inter->errorY + err2Y); - chi2Z = CAMath::Square((float)inter->posZ + deltaZ - posZ) / ((float)inter->errorZ + err2Z); - } else if (mFitInProjections) { - const float Iz0 = inter->posY - mP[0]; - const float Iz1 = inter->posZ + deltaZ - mP[1]; - const float Iw0 = 1.f / (mC[0] + (float)inter->errorY); - const float Iw2 = 1.f / (mC[2] + (float)inter->errorZ); - const float Ik00 = mC[0] * Iw0; - const float Ik11 = mC[2] * Iw2; - const float ImP0 = mP[0] + Ik00 * Iz0; - const float ImP1 = mP[1] + Ik11 * Iz1; - const float ImC0 = mC[0] - Ik00 * mC[0]; - const float ImC2 = mC[2] - Ik11 * mC[2]; - // printf("\t%21sInterpo ----- abde artaf%16s Y %8.3f, Z %8.3f (Errors %f <-- (%f, %f) %f <-- (%f, %f))\n", "", "", ImP0, ImP1, sqrtf(ImC0), sqrtf(mC[0]), sqrtf(inter->errorY), sqrtf(ImC2), sqrtf(mC[2]), sqrtf(inter->errorZ)); - const float Jz0 = posY - ImP0; - const float Jz1 = posZ - ImP1; - const float Jw0 = 1.f / (ImC0 + err2Y); - const float Jw2 = 1.f / (ImC2 + err2Z); - chi2Y = Jw0 * Jz0 * Jz0; - chi2Z = Jw2 * Jz1 * Jz1; - } else { - const float Iz0 = inter->posY - mP[0]; - const float Iz1 = inter->posZ + deltaZ - mP[1]; - float Iw0 = mC[2] + (float)inter->errorZ; - float Iw2 = mC[0] + (float)inter->errorY; - float Idet = CAMath::Max(1e-10f, Iw0 * Iw2 - mC[1] * mC[1]); - Idet = 1.f / Idet; - Iw0 *= Idet; - const float Iw1 = mC[1] * Idet; - Iw2 *= Idet; - const float Ik00 = mC[0] * Iw0 + mC[1] * Iw1; - const float Ik01 = mC[0] * Iw1 + mC[1] * Iw2; - const float Ik10 = mC[1] * Iw0 + mC[2] * Iw1; - const float Ik11 = mC[1] * Iw1 + mC[2] * Iw2; - const float ImP0 = mP[0] + Ik00 * Iz0 + Ik01 * Iz1; - const float ImP1 = mP[1] + Ik10 * Iz0 + Ik11 * Iz1; - const float ImC0 = mC[0] - Ik00 * mC[0] + Ik01 * mC[1]; - const float ImC1 = mC[1] - Ik10 * mC[0] + Ik11 * mC[1]; - const float ImC2 = mC[2] - Ik10 * mC[1] + Ik11 * mC[2]; - const float Jz0 = posY - ImP0; - const float Jz1 = posZ - ImP1; - float Jw0 = ImC2 + err2Z; - float Jw2 = ImC0 + err2Y; - float Jdet = CAMath::Max(1e-10f, Jw0 * Jw2 - ImC1 * ImC1); - Jdet = 1.f / Jdet; - Jw0 *= Jdet; - const float Jw1 = ImC1 * Jdet; - Jw2 *= Jdet; - chi2Y = CAMath::Abs((Jw0 * Jz0 + Jw1 * Jz1) * Jz0); - chi2Z = CAMath::Abs((Jw1 * Jz0 + Jw2 * Jz1) * Jz1); - } - if (RejectCluster(chi2Y * param.rec.tpc.clusterRejectChi2TolleranceY, chi2Z * param.rec.tpc.clusterRejectChi2TolleranceZ, clusterState)) { // TODO: Relative Pt resolution decreases slightly, why? - // printf("Reject Cluster chiy2 %f chiz2 %f (Pos Y: %f - %f %f ; Pos Z: %f - %f %f)\n", chi2Y, chi2Z, posY, mP[0], (float)inter->posY, posZ, mP[1], (float)inter->posZ + deltaZ); - return updateErrorClusterRejectedInInterpolation; - } + inter->posY = mP[0]; + inter->posZ = mP[1]; + if (mT->NDF() <= 0) { + inter->errorY = inter->errorZ = 100.f; + } else { + inter->errorY = mC[0]; + inter->errorZ = mC[2]; + } +} + +GPUd() int32_t GPUTPCGMPropagator::InterpolateReject(const GPUParam& GPUrestrict() param, float posY, float posZ, int16_t clusterState, gputpcgmmergertypes::InterpolationErrorHit* inter, float err2Y, float err2Z, float deltaZ) +{ + float* GPUrestrict() mC = mT->Cov(); + float* GPUrestrict() mP = mT->Par(); + float chi2Y, chi2Z; + if (mT->NDF() <= 0) { + chi2Y = CAMath::Square((float)inter->posY - posY) / ((float)inter->errorY + err2Y); + chi2Z = CAMath::Square((float)inter->posZ + deltaZ - posZ) / ((float)inter->errorZ + err2Z); + } else if (mFitInProjections) { + const float Iz0 = inter->posY - mP[0]; + const float Iz1 = inter->posZ + deltaZ - mP[1]; + const float Iw0 = 1.f / (mC[0] + (float)inter->errorY); + const float Iw2 = 1.f / (mC[2] + (float)inter->errorZ); + const float Ik00 = mC[0] * Iw0; + const float Ik11 = mC[2] * Iw2; + const float ImP0 = mP[0] + Ik00 * Iz0; + const float ImP1 = mP[1] + Ik11 * Iz1; + const float ImC0 = mC[0] - Ik00 * mC[0]; + const float ImC2 = mC[2] - Ik11 * mC[2]; + // printf("\t%21sInterpo ----- abde artaf%16s Y %8.3f, Z %8.3f (Errors %f <-- (%f, %f) %f <-- (%f, %f))\n", "", "", ImP0, ImP1, sqrtf(ImC0), sqrtf(mC[0]), sqrtf(inter->errorY), sqrtf(ImC2), sqrtf(mC[2]), sqrtf(inter->errorZ)); + const float Jz0 = posY - ImP0; + const float Jz1 = posZ - ImP1; + const float Jw0 = 1.f / (ImC0 + err2Y); + const float Jw2 = 1.f / (ImC2 + err2Z); + chi2Y = Jw0 * Jz0 * Jz0; + chi2Z = Jw2 * Jz1 * Jz1; + } else { + const float Iz0 = inter->posY - mP[0]; + const float Iz1 = inter->posZ + deltaZ - mP[1]; + float Iw0 = mC[2] + (float)inter->errorZ; + float Iw2 = mC[0] + (float)inter->errorY; + float Idet = CAMath::Max(1e-10f, Iw0 * Iw2 - mC[1] * mC[1]); + Idet = 1.f / Idet; + Iw0 *= Idet; + const float Iw1 = mC[1] * Idet; + Iw2 *= Idet; + const float Ik00 = mC[0] * Iw0 + mC[1] * Iw1; + const float Ik01 = mC[0] * Iw1 + mC[1] * Iw2; + const float Ik10 = mC[1] * Iw0 + mC[2] * Iw1; + const float Ik11 = mC[1] * Iw1 + mC[2] * Iw2; + const float ImP0 = mP[0] + Ik00 * Iz0 + Ik01 * Iz1; + const float ImP1 = mP[1] + Ik10 * Iz0 + Ik11 * Iz1; + const float ImC0 = mC[0] - Ik00 * mC[0] + Ik01 * mC[1]; + const float ImC1 = mC[1] - Ik10 * mC[0] + Ik11 * mC[1]; + const float ImC2 = mC[2] - Ik10 * mC[1] + Ik11 * mC[2]; + const float Jz0 = posY - ImP0; + const float Jz1 = posZ - ImP1; + float Jw0 = ImC2 + err2Z; + float Jw2 = ImC0 + err2Y; + float Jdet = CAMath::Max(1e-10f, Jw0 * Jw2 - ImC1 * ImC1); + Jdet = 1.f / Jdet; + Jw0 *= Jdet; + const float Jw1 = ImC1 * Jdet; + Jw2 *= Jdet; + chi2Y = CAMath::Abs((Jw0 * Jz0 + Jw1 * Jz1) * Jz0); + chi2Z = CAMath::Abs((Jw1 * Jz0 + Jw2 * Jz1) * Jz1); + } + if (RejectCluster(chi2Y * param.rec.tpc.clusterRejectChi2TolleranceY, chi2Z * param.rec.tpc.clusterRejectChi2TolleranceZ, clusterState)) { // TODO: Relative Pt resolution decreases slightly, why? + // printf("Reject Cluster chiy2 %f chiz2 %f (Pos Y: %f - %f %f ; Pos Z: %f - %f %f)\n", chi2Y, chi2Z, posY, mP[0], (float)inter->posY, posZ, mP[1], (float)inter->posZ + deltaZ); + return updateErrorClusterRejectedInInterpolation; } return 0; } diff --git a/GPU/GPUTracking/Merger/GPUTPCGMPropagator.h b/GPU/GPUTracking/Merger/GPUTPCGMPropagator.h index c3f7aac84a1a5..3f1a74cfb6c26 100644 --- a/GPU/GPUTracking/Merger/GPUTPCGMPropagator.h +++ b/GPU/GPUTracking/Merger/GPUTPCGMPropagator.h @@ -57,11 +57,6 @@ class GPUTPCGMPropagator updateErrorClusterRejectedInUpdate = 4, updateErrorClusterRejectedEdge = 5 }; - enum RejectChi2Mode { - rejectDirect = 1, - rejectInterFill = 2, - rejectInterReject = 3 - }; GPUdDefault() GPUTPCGMPropagator() = default; @@ -104,21 +99,28 @@ class GPUTPCGMPropagator GPUd() int32_t PropagateToXAlphaBz(float posX, float posAlpha, bool inFlyDirection); - GPUd() int32_t Update(float posY, float posZ, int32_t iRow, const GPUParam& param, int16_t clusterState, int8_t rejectChi2, bool refit, int8_t sector, float time, float avgInvCharge, float invCharge); - GPUd() int32_t Update(float posY, float posZ, int32_t iRow, const GPUParam& param, int16_t clusterState, int8_t rejectChi2, bool refit, float err2Y, float err2Z); + GPUd() int32_t Update(float posY, float posZ, int32_t iRow, const GPUParam& param, int16_t clusterState, bool rejectChi2, bool refit, int8_t sector, float time, float avgInvCharge, float invCharge); + GPUd() int32_t Update(float posY, float posZ, int32_t iRow, const GPUParam& param, int16_t clusterState, bool rejectChi2, bool refit, float err2Y, float err2Z); GPUd() int32_t Update(float posY, float posZ, int16_t clusterState, bool rejectChi2, float err2Y, float err2Z, const GPUParam* param = nullptr); - GPUd() int32_t InterpolateReject(const GPUParam& param, float posY, float posZ, int16_t clusterState, int8_t rejectChi2, gputpcgmmergertypes::InterpolationErrorHit* inter, float err2Y, float err2Z, float deltaZ); + GPUd() void InterpolateFill(gputpcgmmergertypes::InterpolationErrorHit* inter); + GPUd() int32_t InterpolateReject(const GPUParam& param, float posY, float posZ, int16_t clusterState, gputpcgmmergertypes::InterpolationErrorHit* inter, float err2Y, float err2Z, float deltaZ); GPUd() float PredictChi2(float posY, float posZ, int32_t iRow, const GPUParam& param, int16_t clusterState, uint8_t sector, float time, float avgCharge, float charge) const; GPUd() float PredictChi2(float posY, float posZ, float err2Y, float err2Z) const; GPUd() static int32_t RejectCluster(float chiY, float chiZ, uint8_t clusterState) { - if (chiY > 9.f || chiZ > 9.f) { // TODO: Check how a track can have chi2/ncl > 18 + if (chiY > 9.f || chiZ > 9.f) { + return 2; + } + if (chiZ > 6.25f && (clusterState & (GPUTPCGMMergedTrackHit::flagSplit | GPUTPCGMMergedTrackHit::flagShared | GPUTPCGMMergedTrackHit::flagEdge | GPUTPCGMMergedTrackHit::flagSingle))) { + return 2; + } + if (chiY > 6.25f && (clusterState & (GPUTPCGMMergedTrackHit::flagSplit | GPUTPCGMMergedTrackHit::flagShared))) { return 2; } - if ((chiY > 6.25f || chiZ > 6.25f) && (clusterState & (GPUTPCGMMergedTrackHit::flagSplit | GPUTPCGMMergedTrackHit::flagShared))) { + if (chiY > 1.f && (clusterState & GPUTPCGMMergedTrackHit::flagEdge)) { return 2; } - if ((chiY > 1.f || chiZ > 6.25f) && (clusterState & (GPUTPCGMMergedTrackHit::flagEdge | GPUTPCGMMergedTrackHit::flagSingle))) { + if (chiY > 2.f && (clusterState & GPUTPCGMMergedTrackHit::flagSingle)) { return 2; } return 0; diff --git a/GPU/GPUTracking/Merger/GPUTPCGMTrackParam.cxx b/GPU/GPUTracking/Merger/GPUTPCGMTrackParam.cxx index 7b70d94b9cfc5..b2851c063a72e 100644 --- a/GPU/GPUTracking/Merger/GPUTPCGMTrackParam.cxx +++ b/GPU/GPUTracking/Merger/GPUTPCGMTrackParam.cxx @@ -13,7 +13,7 @@ /// \author David Rohr, Sergey Gorbunov #define GPUCA_CADEBUG 0 -#define DEBUG_SINGLE_TRACK -1 +// #define DEBUG_REBUILD_MC #include "GPUTPCDef.h" #include "GPUTPCGMTrackParam.h" @@ -36,7 +36,16 @@ #ifdef GPUCA_CADEBUG_ENABLED #include "GPUSettings.h" +#include +#include +#include #include "AliHLTTPCClusterMCData.h" +static thread_local bool CADEBUG_CONDITION = true; +#endif + +#ifndef GPUCA_GPUCODE +#include "SimulationDataFormat/ConstMCTruthContainer.h" +#include "SimulationDataFormat/MCCompLabel.h" #endif #ifndef GPUCA_GPUCODE_DEVICE @@ -47,12 +56,13 @@ using namespace o2::gpu; using namespace o2::tpc; -GPUd() bool GPUTPCGMTrackParam::Fit(GPUTPCGMMerger* GPUrestrict() merger, int32_t iTrk, GPUTPCGMMergedTrackHit* GPUrestrict() clusters, int32_t& GPUrestrict() N, int32_t& GPUrestrict() NTolerated, float& GPUrestrict() Alpha, int32_t attempt, float maxSinPhi, GPUTPCGMMergedTrack& GPUrestrict() track) +GPUd() bool GPUTPCGMTrackParam::Fit(GPUTPCGMMerger& GPUrestrict() merger, int32_t iTrk, int32_t& GPUrestrict() N, int32_t& GPUrestrict() NTolerated, float& GPUrestrict() Alpha, GPUTPCGMMergedTrack& GPUrestrict() track, bool rebuilt, bool retryAttempt) { - constexpr float kDeg2Rad = M_PI / 180.f; - CADEBUG(static constexpr float kSectAngle = 2 * M_PI / 18.f); + constexpr float maxSinPhi = constants::MAX_SIN_PHI; - const GPUParam& GPUrestrict() param = merger->Param(); + const GPUParam& GPUrestrict() param = merger.Param(); + const float maxSinForUpdate = param.rec.tpc.maxSinPhiForFitUpdate; + GPUTPCGMMergedTrackHit* GPUrestrict() clusters = merger.Clusters() + track.FirstClusterRef(); GPUdEdx dEdx, dEdxAlt; GPUTPCGMPropagator prop; @@ -60,63 +70,69 @@ GPUd() bool GPUTPCGMTrackParam::Fit(GPUTPCGMMerger* GPUrestrict() merger, int32_ prop.SetMaterialTPC(); prop.SetPolynomialField(¶m.polynomialField); prop.SetMaxSinPhi(maxSinPhi); - if (param.rec.tpc.mergerInterpolateErrors) { - for (int32_t i = 0; i < N; i++) { - interpolation.hit[i].errorY = -1; + if (param.rec.tpc.mergerInterpolateErrors && !rebuilt) { + for (uint32_t i = 0; i < interpolation.size; i++) { // TODO: Tune the zeroing range + interpolation.hit[i].markInvalid(); + } + if (param.rec.tpc.rebuildTrackInFit) { + merger.TrackRebuildHelper()[iTrk].highInclRowLow = merger.TrackRebuildHelper()[iTrk].highInclRowHigh = 255; } } const int32_t nWays = param.rec.tpc.nWays; const int32_t maxN = N; + if (rebuilt && maxN > 0 && param.par.continuousTracking && !track.CCE() && !track.Looper()) { // The rebuilt track can have other first / last clusters than the track the Z shift was computed for + const float t0 = merger.GetConstantMem()->ioPtrs.clustersNative->clustersLinear[clusters[0].num].getTime(), tN = merger.GetConstantMem()->ioPtrs.clustersNative->clustersLinear[clusters[maxN - 1].num].getTime(); + ConstrainZ(merger, clusters[0].sector, CAMath::Max(t0, tN), CAMath::Min(t0, tN)); + } int32_t ihitStart = 0; + int32_t interpolatedStart = 0; float covYYUpd = 0.f; - float lastUpdateX = -1.f; - uint8_t lastRow = 255; - uint8_t lastSector = 255; float deltaZ = 0.f; - for (int32_t iWay = 0; iWay < nWays; iWay++) { + // TODO: Could do 5 ways and interpolate again on the rebuilt clusters and reject in the final way based on that (and get the outer param from the rebuilt clusters), but costs 2 more ways + for (int32_t iWay = rebuilt ? nWays - 1 : 0; iWay < nWays; iWay++) { // TODO DR: Unrolling has no performance improvement on GPU, why? int32_t nMissed = 0, nMissed2 = 0; - float sumInvSqrtCharge = 0.f; + float sumInvSqrtCharge = 0.f; // TODO: Compute in first iteration and store! int32_t nAvgCharge = 0; if (iWay && (iWay & 1) == 0) { - StoreOuter(&track.OuterParam(), prop.GetAlpha()); + StoreOuter(&track.OuterParam(), prop.GetAlpha()); // TODO: Fix outer param! } - int32_t resetT0 = initResetT0(); + int32_t resetT0 = CAMath::Max(10, CAMath::Min(40, 150.f / CAMath::Abs(mP[4]))); const bool refit = (nWays == 1 || iWay >= 1); const bool finalOutInFit = iWay + 2 >= nWays; const bool finalFit = iWay == nWays - 1; - const float maxSinForUpdate = CAMath::Sin(70.f * kDeg2Rad); ResetCovariance(); - prop.SetSeedingErrors(!(refit && attempt == 0)); + prop.SetSeedingErrors(!(refit && !retryAttempt)); prop.SetFitInProjections(true); // param.rec.fitInProjections == -1 ? (iWay == 0) : param.rec.fitInProjections); // TODO: Reenable once fixed prop.SetPropagateBzOnly(param.rec.fitPropagateBzOnly == -1 ? !finalFit : param.rec.fitPropagateBzOnly); - prop.SetMatLUT((param.rec.useMatLUT && finalFit) ? merger->GetConstantMem()->calibObjects.matLUT : nullptr); - prop.SetTrack(this, iWay ? prop.GetAlpha() : Alpha); + prop.SetMatLUT((param.rec.useMatLUT && finalFit) ? merger.GetConstantMem()->calibObjects.matLUT : nullptr); + prop.SetTrack(this, iWay && !rebuilt ? prop.GetAlpha() : Alpha); ConstrainSinPhi(iWay == 0 ? 0.95f : constants::MAX_SIN_PHI_LOW); - CADEBUG(printf("Fitting track %d way %d (sector %d, alpha %f) !!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!\n", iTrk, iWay, CAMath::Float2IntRn(prop.GetAlpha() / kSectAngle) + (mP[1] < 0 ? 18 : 0), prop.GetAlpha())); + CADEBUG(printf("Fitting track %d way %d (sector %d, alpha %f) !!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!\n", iTrk, iWay, CAMath::Float2IntRn(prop.GetAlpha() / GPUTPCGeometry::kSectAngle()) + (mP[1] < 0 ? 18 : 0), prop.GetAlpha())); N = 0; - lastUpdateX = -1; + uint8_t lastUpdateRow = 255, lastPropagateRow = 255, lastSector = 255; + bool fitBroken = false; + float lastUpdateX = -1; const bool inFlyDirection = iWay & 1; const int32_t wayDirection = (iWay & 1) ? -1 : 1; - int32_t goodRows = 0; - for (int32_t ihit = ihitStart; ihit >= 0 && ihit < maxN; ihit += wayDirection) { - const bool crossCE = lastSector != 255 && ((lastSector < 18) ^ (clusters[ihit].sector < 18)); - if (crossCE) { - lastSector = clusters[ihit].sector; - } - - if ((param.rec.tpc.trackFitRejectMode > 0 && nMissed >= param.rec.tpc.trackFitRejectMode) || nMissed2 >= param.rec.tpc.trackFitMaxRowMissedHard || clusters[ihit].state & GPUTPCGMMergedTrackHit::flagReject) { - CADEBUG(printf("\tSkipping hit %d, %d hits rejected, flag %X\n", ihit, nMissed, (int32_t)clusters[ihit].state)); - if (finalOutInFit && !(clusters[ihit].state & GPUTPCGMMergedTrackHit::flagReject)) { - clusters[ihit].state |= GPUTPCGMMergedTrackHit::flagRejectErr; + for (int32_t ihit = ihitStart, interpolationIndex = interpolatedStart; ihit >= 0 && ihit < maxN; ihit += wayDirection, interpolationIndex += wayDirection) { + if (!param.rec.tpc.rebuildTrackInFit || rebuilt) { + if ((param.rec.tpc.trackFitRejectMode > 0 && nMissed >= param.rec.tpc.trackFitRejectMode) || nMissed2 >= param.rec.tpc.trackFitMaxRowMissedHard || (clusters[ihit].state & GPUTPCGMMergedTrackHit::flagReject) || (rebuilt && (clusters[ihit].state & GPUTPCGMMergedTrackHit::flagHighIncl))) { + CADEBUG(printf("\tSkipping hit %d, %d hits rejected, flag %X\n", ihit, nMissed, (int32_t)clusters[ihit].state)); + if (rebuilt && (clusters[ihit].state & GPUTPCGMMergedTrackHit::flagHighIncl)) { + NTolerated++; + } + if (finalOutInFit && !(clusters[ihit].state & (GPUTPCGMMergedTrackHit::flagReject | GPUTPCGMMergedTrackHit::flagHighIncl))) { + clusters[ihit].state |= GPUTPCGMMergedTrackHit::flagRejectErr; + } + continue; } - continue; } const bool allowChangeClusters = finalOutInFit && (nWays == 1 || ((iWay & 1) ? (ihit <= CAMath::Max(maxN / 2, maxN - 30)) : (ihit >= CAMath::Min(maxN / 2, 30)))); @@ -125,217 +141,353 @@ GPUd() bool GPUTPCGMTrackParam::Fit(GPUTPCGMMerger* GPUrestrict() merger, int32_ uint8_t clusterState = clusters[ihit].state; const float clAlpha = param.Alpha(clusters[ihit].sector); float xx, yy, zz; - { - const ClusterNative& GPUrestrict() cl = merger->GetConstantMem()->ioPtrs.clustersNative->clustersLinear[clusters[ihit].num]; - merger->GetConstantMem()->calibObjects.fastTransform->Transform(clusters[ihit].sector, clusters[ihit].row, cl.getPad(), cl.getTime(), xx, yy, zz, mTOffset); - } - // clang-format off + const int32_t currentClusterStatus = MergeDoubleRowClusters(ihit, wayDirection, clusters, merger, prop, xx, yy, zz, maxN, clAlpha, clusterState, param.rec.tpc.rebuildTrackInFit ? rebuilt : allowChangeClusters); + // TODO: Check about tracks who have clusters in the same row multiple times in different sectors + + const auto& cluster = clusters[ihit]; + CADEBUG(printf("\tSector %2d %11sTrack Alpha %8.3f %s, X %8.3f - Y %8.3f, Z %8.3f - QPt %7.2f (%7.2f), SP %5.2f (%5.2f) %28s --- Cov sY %8.3f sZ %8.3f sSP %8.3f sPt %8.3f - YPt %8.3f\n", (int32_t)cluster.sector, "", prop.GetAlpha(), (CAMath::Abs(prop.GetAlpha() - clAlpha) < 0.01 ? " " : " R!"), mX, mP[0], mP[1], mP[4], prop.GetQPt0(), mP[2], prop.GetSinPhi0(), "", sqrtf(mC[0]), sqrtf(mC[2]), sqrtf(mC[5]), sqrtf(mC[14]), mC[10])); CADEBUG(printf("\tHit %3d/%3d Row %3d: Cluster Alpha %8.3f %3d, X %8.3f - Y %8.3f, Z %8.3f (Missed %d)\n", ihit, maxN, (int32_t)clusters[ihit].row, clAlpha, (int32_t)clusters[ihit].sector, xx, yy, zz, nMissed)); - // CADEBUG(if ((uint32_t)merger->GetTrackingChain()->mIOPtrs.nMCLabelsTPC > clusters[ihit].num)) - // CADEBUG({printf(" MC:"); for (int32_t i = 0; i < 3; i++) {int32_t mcId = merger->GetTrackingChain()->mIOPtrs.mcLabelsTPC[clusters[ihit].num].fClusterID[i].fMCID; if (mcId >= 0) printf(" %d", mcId); } } printf("\n")); - // clang-format on - if (MergeDoubleRowClusters(ihit, wayDirection, clusters, merger, prop, xx, yy, zz, maxN, clAlpha, clusterState, allowChangeClusters) == -1) { - nMissed++; - nMissed2++; - continue; - } - if (param.rec.tpc.rejectIFCLowRadiusCluster) { - const float r2 = xx * xx + yy * yy; - const float rmax = (83.5f + param.rec.tpc.sysClusErrorMinDist); - if (r2 < rmax * rmax) { - MarkClusters(clusters, ihitMergeFirst, ihit, wayDirection, GPUTPCGMMergedTrackHit::flagRejectErr); - continue; - } + const int32_t rowGap = lastPropagateRow == 255 ? 0 : CAMath::Max(0, CAMath::Abs(cluster.row - lastPropagateRow) - 1); + if ((uint32_t)(interpolationIndex + rowGap * wayDirection) >= interpolation.size) { // covers all indices of the row gap, since they are monotonic + merger.raiseError(GPUErrors::ERROR_MERGER_INTERPOLATION_OVERFLOW, interpolationIndex + rowGap * wayDirection, interpolation.size); + fitBroken = true; + break; } - const auto& cluster = clusters[ihit]; - - // clang-format off - CADEBUG(printf("\tSector %2d %11sTrack Alpha %8.3f %s, X %8.3f - Y %8.3f, Z %8.3f - QPt %7.2f (%7.2f), SP %5.2f (%5.2f) %28s --- Cov sY %8.3f sZ %8.3f sSP %8.3f sPt %8.3f - YPt %8.3f\n", (int32_t)cluster.sector, "", prop.GetAlpha(), (CAMath::Abs(prop.GetAlpha() - clAlpha) < 0.01 ? " " : " R!"), mX, mP[0], mP[1], mP[4], prop.GetQPt0(), mP[2], prop.GetSinPhi0(), "", sqrtf(mC[0]), sqrtf(mC[2]), sqrtf(mC[5]), sqrtf(mC[14]), mC[10])); - // clang-format on - if (allowChangeClusters && lastRow != 255 && CAMath::Abs(cluster.row - lastRow) > 1) { - if GPUCA_RTC_CONSTEXPR (GPUCA_GET_CONSTEXPR(param.par, dodEdx)) { - bool dodEdx = param.dodEdxEnabled && param.rec.tpc.adddEdxSubThresholdClusters && finalFit && CAMath::Abs(cluster.row - lastRow) == 2; - dodEdx = AttachClustersPropagate(merger, cluster.sector, lastRow, cluster.row, iTrk, track.Leg() == 0, prop, inFlyDirection, constants::MAX_SIN_PHI, dodEdx); - if (dodEdx) { - dEdx.fillSubThreshold(lastRow - wayDirection); - if GPUCA_RTC_CONSTEXPR (GPUCA_GET_CONSTEXPR(param.rec.tpc, dEdxClusterRejectionFlagMask) != GPUCA_GET_CONSTEXPR(param.rec.tpc, dEdxClusterRejectionFlagMaskAlt)) { - dEdxAlt.fillSubThreshold(lastRow - wayDirection); + uint8_t dEdxSubThresholdRow = 255; + if (rowGap) { + const bool dodEdx = param.dodEdxEnabled && param.rec.tpc.adddEdxSubThresholdClusters && finalFit && CAMath::Abs(cluster.row - lastUpdateRow) == 2 && cluster.sector == lastSector && currentClusterStatus == 0; + const bool doAttach = allowChangeClusters && !param.rec.tpc.rebuildTrackInFit && !(merger.Param().rec.tpc.disableMarkAdjacent & 2); + const bool crossCE = lastSector != 255 && ((cluster.sector >= GPUTPCGeometry::NSECTORS / 2) != (lastSector >= GPUTPCGeometry::NSECTORS / 2)); + const bool ceOverlap = crossCE && ((cluster.row > lastPropagateRow) != ((clusters[0].row < clusters[maxN - 1].row) ^ (iWay & 1))); // Row gap goes backwards at the CE crossing, these rows were already searched on both sides + const uint8_t doInterpolate = ((param.rec.tpc.disableRebuildAttachment & 2) || ceOverlap) ? 0 : ((param.rec.tpc.rebuildTrackInFit && iWay == nWays - 3) ? 1 : ((param.rec.tpc.rebuildTrackInFit && iWay == nWays - 2) ? 2 : 0)); + if (lastUpdateRow != 255 && (dodEdx || doAttach || doInterpolate)) { + int32_t step = cluster.row > lastPropagateRow ? 1 : -1; + uint8_t sector = lastSector; + + for (int32_t iRow = lastPropagateRow + step, index = interpolationIndex; iRow != cluster.row; iRow += step, index += wayDirection) { + float tmpX, tmpY, tmpZ; + if (prop.GetPropagatedYZ(GPUTPCGeometry::Row2X(iRow), tmpY, tmpZ)) { + break; + } + merger.GetConstantMem()->calibObjects.fastTransform->InverseTransformYZtoX(sector, iRow, tmpY, tmpZ, tmpX); + if (prop.PropagateToXAlpha(tmpX, prop.GetAlpha(), inFlyDirection)) { + break; + } + FitAddRow(iRow, sector, iTrk, track, prop, inFlyDirection, merger, &dEdxSubThresholdRow, dodEdx, doAttach, doInterpolate, interpolation.hit[index], deltaZ, sumInvSqrtCharge, nAvgCharge); + if (crossCE && doInterpolate == 2) { // Track crosses the CE within the row gap, search also on the other side, the sector there has the same alpha + FitAddRow(iRow, sector >= GPUTPCGeometry::NSECTORS / 2 ? sector - GPUTPCGeometry::NSECTORS / 2 : sector + GPUTPCGeometry::NSECTORS / 2, iTrk, track, prop, inFlyDirection, merger, &dEdxSubThresholdRow, false, false, doInterpolate, interpolation.hit[index], deltaZ, sumInvSqrtCharge, nAvgCharge); } } } + interpolationIndex += rowGap * wayDirection; } + lastPropagateRow = cluster.row; + auto& inter = interpolation.hit[interpolationIndex]; int32_t retValProp = prop.PropagateToXAlpha(xx, clAlpha, inFlyDirection); - // clang-format off - CADEBUG(if (!CheckCov()){printf("INVALID COV AFTER PROPAGATE!!!\n");}); - // clang-format on - if (retValProp == -2) // Rotation failed, try to bring to new x with old alpha first, rotate, and then propagate to x, alpha - { - CADEBUG(printf("REROTATE\n")); - if (prop.PropagateToXAlpha(xx, prop.GetAlpha(), inFlyDirection) == 0) { - retValProp = prop.PropagateToXAlpha(xx, clAlpha, inFlyDirection); + if ((retValProp == -2 && // Rotation failed, try to bring to new x with old alpha first, rotate, and then propagate to x, alpha + (prop.PropagateToXAlpha(xx, prop.GetAlpha(), inFlyDirection) != 0 || // Cannot rotate to new alpha at all + prop.PropagateToXAlpha(xx, clAlpha, inFlyDirection))) || // propagation fails nonetheless + retValProp) { // failed for other reason but rotation + CADEBUG(printf(" --- break-prop\n")); + MarkClusters(clusters, ihitMergeFirst, ihit, wayDirection, GPUTPCGMMergedTrackHit::flagHighIncl); + nMissed2++; + NTolerated++; + if (param.rec.tpc.rebuildTrackInFit && iWay == nWays - 2 && param.rec.tpc.rebuildTrackMaxNonIntCov > 0 && !(param.rec.tpc.disableRebuildAttachment & 4)) { + FindBestInterpolatedHit(merger, inter, cluster.sector, cluster.row, deltaZ, sumInvSqrtCharge, nAvgCharge, prop, iTrk, true); // TODO: mark clusters from single-sided finding as dubious, and apply stricter restriction cut later } - } - if (lastRow == 255 || CAMath::Abs((int32_t)lastRow - (int32_t)cluster.row) > 5 || lastSector != cluster.sector || (param.rec.tpc.trackFitRejectMode < 0 && -nMissed <= param.rec.tpc.trackFitRejectMode)) { - goodRows = 0; - } else { - goodRows++; - } - if (retValProp == 0) { - lastRow = cluster.row; - lastSector = cluster.sector; + continue; } // clang-format off - 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 - PErr %d", "", prop.GetAlpha(), mX, mP[0], mP[1], mP[4], prop.GetQPt0(), mP[2], prop.GetSinPhi0(), mP[0] - yy, mP[1] - zz, sqrtf(mC[0]), sqrtf(mC[2]), sqrtf(mC[5]), sqrtf(mC[14]), mC[10], retValProp)); + CADEBUG(if (!CheckCov()){printf("INVALID COV AFTER PROPAGATE!!!\n");}); + 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 - PErr %d\n", "", prop.GetAlpha(), mX, mP[0], mP[1], mP[4], prop.GetQPt0(), mP[2], prop.GetSinPhi0(), mP[0] - yy, mP[1] - zz, sqrtf(mC[0]), sqrtf(mC[2]), sqrtf(mC[5]), sqrtf(mC[14]), mC[10], retValProp)); // clang-format on - - if (crossCE) { - if (param.rec.tpc.addErrorsCECrossing) { - if (param.rec.tpc.addErrorsCECrossing >= 2) { - AddCovDiagErrorsWithCorrelations(param.rec.tpc.errorsCECrossing); - } else { - AddCovDiagErrors(param.rec.tpc.errorsCECrossing); - } - } else if (mC[2] < 0.5f) { - mC[2] = 0.5f; + if (mNDF >= 0 && (mC[0] > param.rec.tpc.trackFitCovLimit || mC[2] > param.rec.tpc.trackFitCovLimit)) { + if (param.rec.tpc.rebuildTrackInFit && iWay == nWays - 2 && param.rec.tpc.rebuildTrackMaxNonIntCov > 0 && !(param.rec.tpc.disableRebuildAttachment & 4)) { + InterpolateMissingRows(merger, interpolation, clusters, ihit, interpolationIndex, cluster.row, deltaZ, sumInvSqrtCharge, nAvgCharge, prop, iTrk); } + fitBroken = true; + break; // bad chi2 for the whole track, stop the fit } float uncorrectedY = -1e6f; - if (allowChangeClusters) { + if (param.rec.tpc.rebuildTrackInFit) { + if (iWay == nWays - 2 && !(param.rec.tpc.disableRebuildAttachment & 1)) { + uncorrectedY = FindBestInterpolatedHit(merger, inter, cluster.sector, cluster.row, deltaZ, sumInvSqrtCharge, nAvgCharge, prop, iTrk, false); // TODO: Check whether we do not need to use track Y for some rejection steps + } + if (allowChangeClusters) { + AttachClusters(merger, cluster.sector, cluster.row, iTrk, track.Leg() == 0, prop); // TODO: Do this during FindBestInterpolatedHit + } + } else if (allowChangeClusters) { uncorrectedY = AttachClusters(merger, cluster.sector, cluster.row, iTrk, track.Leg() == 0, prop); } + if (param.rec.tpc.rejectEdgeClustersInTrackFit && uncorrectedY <= -1e6f) { + float tmpZ; + merger.GetConstantMem()->calibObjects.fastTransform->InverseTransformYZtoNominalYZ(cluster.sector, cluster.row, mP[0], mP[1], uncorrectedY, tmpZ); + } - const bool sinPhiErr = mNDF > 0 && CAMath::Abs(prop.GetSinPhi0()) >= maxSinForUpdate; - if (mNDF >= 0 && (mC[0] > param.rec.tpc.trackFitCovLimit || mC[2] > param.rec.tpc.trackFitCovLimit)) { - break; + HandleCrossCE(param, cluster.sector, lastSector); + lastSector = cluster.sector; + + if (param.rec.tpc.mergerInterpolateErrors && iWay == nWays - 3) { + prop.InterpolateFill(&inter); } - if (retValProp || sinPhiErr) { - MarkClusters(clusters, ihitMergeFirst, ihit, wayDirection, GPUTPCGMMergedTrackHit::flagHighIncl); + + if (currentClusterStatus) { + nMissed++; nMissed2++; - NTolerated++; - CADEBUG(printf(", %d --- break\n", (int32_t)sinPhiErr)); continue; } - CADEBUG(printf("\n")); - int32_t retValUpd = 0, retValInt = 0; - float threshold = 3.f + (lastUpdateX >= 0 ? (CAMath::Abs(mX - lastUpdateX) / 2) : 0.f); - if (mNDF > (int32_t)param.rec.tpc.mergerNonInterpolateRejectMinNDF && (CAMath::Abs(yy - mP[0]) > threshold || CAMath::Abs(zz - mP[1]) > threshold)) { - retValUpd = GPUTPCGMPropagator::updateErrorClusterRejectedDistance; - } else { - int8_t rejectChi2 = 0; - if (attempt == 0) { - if (param.rec.tpc.mergerInterpolateErrors && CAMath::Abs(ihit - ihitMergeFirst) <= 1) { - if (iWay == nWays - 3) { - rejectChi2 = GPUTPCGMPropagator::rejectInterFill; - } else if (iWay == nWays - 2) { - rejectChi2 = GPUTPCGMPropagator::rejectInterReject; - } else if (iWay == nWays - 1) { - rejectChi2 = (param.rec.tpc.mergerInterpolateRejectAlsoOnCurrentPosition && GetNDF() > (int32_t)param.rec.tpc.mergerNonInterpolateRejectMinNDF) ? GPUTPCGMPropagator::rejectDirect : 0; - } + if (mNDF > 0 && CAMath::Abs(prop.GetSinPhi0()) >= maxSinForUpdate) { // TODO: If NDF is large enough, and mP[2] is not yet out of range, reinit linearization + const bool inward = clusters[0].row > clusters[maxN - 1].row; + const bool markHighIncl = (mP[2] > 0) ^ (mP[4] < 0) ^ inward ^ (iWay & 1); + if (mNDF > 10 && CAMath::Abs(mP[2]) < maxSinForUpdate && markHighIncl) { + CADEBUG(printf("Reinit linearization\n")); + prop.SetTrack(this, prop.GetAlpha()); + } else { + if (!markHighIncl && nWays >= 2 && CAMath::Abs(mP[2]) < maxSinPhi && CAMath::Max(CAMath::Abs(yy - mP[0]), CAMath::Abs(zz - mP[1])) <= 5.f) { // abs(sinPhi) decreases in the direction we are going (we come from another sector) and the cluster is close: continue from the cluster position, a normal update could distort sinPhi and QPt at this inclination + float err2Y, err2Z; + prop.GetErr2(err2Y, err2Z, param, zz, cluster.row, clusterState, cluster.sector, -1.f, 0.f, 0.f); + prop.Model().Y() = mP[0] = yy; // Like an update that takes the position fully from the cluster and leaves the other parameters untouched + prop.Model().Z() = mP[1] = zz; + mC[0] = err2Y; + mC[2] = err2Z; + mC[1] = mC[3] = mC[4] = mC[6] = mC[7] = mC[10] = mC[11] = 0.f; + lastUpdateX = mX; } else { - rejectChi2 = allowChangeClusters && goodRows > 5; + nMissed2++; } - } - - float err2Y, err2Z; - const float time = merger->GetConstantMem()->ioPtrs.clustersNative ? merger->GetConstantMem()->ioPtrs.clustersNative->clustersLinear[cluster.num].getTime() : -1.f; - const float invSqrtCharge = merger->GetConstantMem()->ioPtrs.clustersNative ? CAMath::InvSqrt(merger->GetConstantMem()->ioPtrs.clustersNative->clustersLinear[cluster.num].qMax) : 0.f; - const float invCharge = merger->GetConstantMem()->ioPtrs.clustersNative ? (1.f / merger->GetConstantMem()->ioPtrs.clustersNative->clustersLinear[cluster.num].qMax) : 0.f; - float invAvgCharge = (sumInvSqrtCharge += invSqrtCharge) / ++nAvgCharge; - invAvgCharge *= invAvgCharge; - prop.GetErr2(err2Y, err2Z, param, zz, cluster.row, clusterState, cluster.sector, time, invAvgCharge, invCharge); - - if (rejectChi2 >= GPUTPCGMPropagator::rejectInterFill) { - if (rejectChi2 == GPUTPCGMPropagator::rejectInterReject && interpolation.hit[ihit].errorY < (GPUCA_PAR_MERGER_INTERPOLATION_ERROR_TYPE_A)0) { - rejectChi2 = GPUTPCGMPropagator::rejectDirect; - } else { - retValInt = prop.InterpolateReject(param, yy, zz, clusterState, rejectChi2, &interpolation.hit[ihit], err2Y, err2Z, deltaZ); + MarkClusters(clusters, ihitMergeFirst, ihit, wayDirection, GPUTPCGMMergedTrackHit::flagHighIncl); + CADEBUG(printf(" --- break-sinphi\n")); + NTolerated++; + if (param.rec.tpc.rebuildTrackInFit && markHighIncl) { + if (inward ^ (iWay & 1)) { + if (merger.TrackRebuildHelper()[iTrk].highInclRowLow == 255) { + merger.TrackRebuildHelper()[iTrk].highInclRowLow = cluster.row; + } + } else { + if (merger.TrackRebuildHelper()[iTrk].highInclRowHigh == 255) { + merger.TrackRebuildHelper()[iTrk].highInclRowHigh = cluster.row; + } + } + // TODO: We can perhaps break here, if we pick up remaining rows } + continue; } + } - if (param.rec.tpc.rejectEdgeClustersInTrackFit && uncorrectedY > -1e6f && param.rejectEdgeClusterByY(uncorrectedY, cluster.row, CAMath::Sqrt(mC[0]))) { // uncorrectedY > -1e6f implies allowChangeClusters - retValUpd = GPUTPCGMPropagator::updateErrorClusterRejectedEdge; - } else { - retValUpd = prop.Update(yy, zz, cluster.row, param, clusterState, rejectChi2, refit, err2Y, err2Z); + int32_t retValHit = FitHit(merger, iTrk, track, xx, yy, zz, clusterState, clAlpha, iWay, inFlyDirection, deltaZ, lastUpdateX, clusters, prop, inter, dEdx, dEdxAlt, sumInvSqrtCharge, nAvgCharge, ihit, ihitMergeFirst, allowChangeClusters, refit, finalFit, nMissed, nMissed2, uncorrectedY, retryAttempt); + if (retValHit == 0) { + DodEdx(dEdx, dEdxAlt, merger, finalFit, ihit, ihitMergeFirst, wayDirection, clusters, clusterState, zz, dEdxSubThresholdRow GPUCA_DEBUG_STREAMER_CHECK(, iTrk)); + if (CAMath::Abs(mP[4]) * param.qptB5Scaler > 10 && --resetT0 <= 0 && CAMath::Abs(mP[2]) < 0.15f && CAMath::Square(mP[0] - prop.Model().Y()) + CAMath::Square(mP[1] - prop.Model().Z()) > 1) { // TODO: Could be investigated to reinit the linearization also when sinPhi drifts away from it (a reinit when QPt differs makes things worse, early in the way the QPt of the linearization is usually the better one) + CADEBUG(printf("Reinit linearization: way %d, QPt %7.2f (%7.2f), SP %5.2f (%5.2f), NDF %d\n", iWay, mP[4], prop.GetQPt0(), mP[2], prop.GetSinPhi0(), mNDF)); + prop.SetTrack(this, prop.GetAlpha()); + resetT0 = 20; } - GPUCA_DEBUG_STREAMER_CHECK(if (o2::utils::DebugStreamer::checkStream(o2::utils::StreamFlags::streamUpdateTrack, iTrk)) { - merger->DebugStreamerUpdate(iTrk, ihit, xx, yy, zz, cluster, merger->GetConstantMem()->ioPtrs.clustersNative->clustersLinear[cluster.num], *this, prop, interpolation.hit[ihit], rejectChi2, refit, retValUpd, sumInvSqrtCharge / nAvgCharge * sumInvSqrtCharge / nAvgCharge, yy, zz, clusterState, retValInt, err2Y, err2Z); - }); - } - // clang-format off - CADEBUG(if (!CheckCov()) GPUError("INVALID COV AFTER UPDATE!!!")); - CADEBUG(printf("\t%21sFit Alpha %8.3f , X %8.3f - Y %8.3f, Z %8.3f - QPt %7.2f (%7.2f), SP %5.2f (%5.2f), DzDs %5.2f %16s --- Cov sY %8.3f sZ %8.3f sSP %8.3f sPt %8.3f - YPt %8.3f - FErr %d %d\n", "", prop.GetAlpha(), mX, mP[0], mP[1], mP[4], prop.GetQPt0(), mP[2], prop.GetSinPhi0(), mP[3], "", sqrtf(mC[0]), sqrtf(mC[2]), sqrtf(mC[5]), sqrtf(mC[14]), mC[10], retValUpd, retValInt)); - // clang-format on - ConstrainSinPhi(); // TODO: Limit using ConstrainSinPhi everywhere! - if (!retValUpd && !retValInt) // track is updated - { - lastUpdateX = mX; - covYYUpd = mC[0]; - nMissed = nMissed2 = 0; - UnmarkClusters(clusters, ihitMergeFirst, ihit, wayDirection, GPUTPCGMMergedTrackHit::flagHighIncl); - N++; ihitStart = ihit; - float dy = mP[0] - prop.Model().Y(); - float dz = mP[1] - prop.Model().Z(); - if (CAMath::Abs(mP[4]) * param.qptB5Scaler > 10 && --resetT0 <= 0 && CAMath::Abs(mP[2]) < 0.15f && dy * dy + dz * dz > 1) { - CADEBUG(printf("Reinit linearization\n")); - prop.SetTrack(this, prop.GetAlpha()); + interpolatedStart = interpolationIndex; + N++; + covYYUpd = mC[0]; + if (param.rec.tpc.rebuildTrackInFit && !rebuilt) { // The high inclination row range is valid only if it lasts until the end of the way, reset it since this later hit was fitted + auto& highInclRow = ((clusters[0].row > clusters[maxN - 1].row) ^ (iWay & 1)) ? merger.TrackRebuildHelper()[iTrk].highInclRowLow : merger.TrackRebuildHelper()[iTrk].highInclRowHigh; + if (highInclRow != 255) { + highInclRow = 255; + } } - if GPUCA_RTC_CONSTEXPR (GPUCA_GET_CONSTEXPR(param.par, dodEdx)) { - if (param.dodEdxEnabled && finalFit) { // TODO: Costimize flag to remove, and option to remove double-clusters - bool acc = (clusterState & param.rec.tpc.dEdxClusterRejectionFlagMask) == 0, accAlt = (clusterState & param.rec.tpc.dEdxClusterRejectionFlagMaskAlt) == 0; - if (acc || accAlt) { - float qtot = 0, qmax = 0, pad = 0, relTime = 0; - const int32_t clusterCount = (ihit - ihitMergeFirst) * wayDirection + 1; - for (int32_t iTmp = ihitMergeFirst; iTmp != ihit + wayDirection; iTmp += wayDirection) { - const ClusterNative& cl = merger->GetConstantMem()->ioPtrs.clustersNative->clustersLinear[cluster.num]; - qtot += cl.getQtot(); - qmax = CAMath::Max(qmax, cl.qMax); - pad += cl.getPad(); - relTime += cl.getTime(); - } - qtot /= clusterCount; // TODO: Weighted Average - pad /= clusterCount; - relTime /= clusterCount; - relTime = relTime - CAMath::Round(relTime); - if (acc) { - dEdx.fillCluster(qtot, qmax, cluster.row, cluster.sector, mP[2], mP[3], merger->GetConstantMem()->calibObjects, zz, pad, relTime GPUCA_DEBUG_STREAMER_CHECK(, iTrk, clusterState)); + } else if (retValHit == 1) { + if (param.rec.tpc.rebuildTrackInFit && iWay == nWays - 2 && param.rec.tpc.rebuildTrackMaxNonIntCov > 0 && !(param.rec.tpc.disableRebuildAttachment & 4)) { + InterpolateMissingRows(merger, interpolation, clusters, ihit + wayDirection, interpolationIndex + wayDirection, cluster.row, deltaZ, sumInvSqrtCharge, nAvgCharge, prop, iTrk); + } + fitBroken = true; + break; + } + + lastUpdateRow = cluster.row; + assert(!param.rec.tpc.mergerInterpolateErrors || rebuilt || iWay != nWays - 2 || ihit || interpolationIndex == 0); + } + if (!(iWay & 1) && !finalFit && !track.CCE() && !track.Looper()) { + deltaZ = ShiftZ(clusters, merger, maxN); + } else { + deltaZ = 0.f; + } + + if (param.rec.tpc.rebuildTrackInFit && !rebuilt && !(param.rec.tpc.disableRebuildAttachment & 16) && param.rec.tpc.rebuildTrackMaxNonIntCov > 0 && iWay >= nWays - 3 && !(fitBroken && iWay == nWays - 2) && CAMath::Abs(mP[2]) < maxSinForUpdate && lastUpdateRow != 255 && mNDF > (int32_t)param.rec.tpc.rebuildTrackNonIntMinNDF) { + const int32_t up = ((clusters[0].row < clusters[maxN - 1].row) ^ (iWay & 1)) ? 1 : -1; + int32_t sector = lastSector; + CADEBUG(if (merger.MergedTrackMC()) printf("Extrapolate Start Track %d - sector %2d row %3d %s - fake %d\n", iTrk, sector, (int32_t)lastPropagateRow, up == 1 ? "upwards" : "downwards", (int)merger.MergedTrackMC()[iTrk].isFake())); + uint8_t rowGapActive = 0, rowGapTotal = 0, missingRowsTotal = 0; + uint8_t lastGoodRow = lastPropagateRow, lastExtrapolateRow = lastPropagateRow, lastFoundRow = lastPropagateRow; + uint8_t consecGoodRows = ihitStart == ((iWay & 1) ? 0 : maxN - 1) ? param.rec.tpc.rebuildTrackExtrMinConsecGoodRows : 0, consecGoodRowsMissing = 0, segmentRows = 0; // Continue the track directly only if its last hit was accepted in the fit + float segmentStartChi2 = mChi2; + const float qPtStart = mP[4], alphaStart = prop.GetAlpha(); + const GPUTPCGMTrackParam paramStart = *this; // To revert to if the extrapolation fails + float lastFoundX = 0.f, lastFoundAlpha = 0.f; + bool revert = false; + prop.SetTrack(this, prop.GetAlpha()); + uint32_t iExtrapolatedRow = lastPropagateRow + up; + for (; iExtrapolatedRow < GPUTPCGeometry::NROWS; iExtrapolatedRow += up) { // uint8_t implies > 0! // TODO: Try to reduce some variables to int8/uint8 to save registers + bool fail = false; + for (int32_t iAttempt = 0; iAttempt < 2; iAttempt++) { + float tmpX, tmpY, tmpZ; + if (prop.GetPropagatedYZ(GPUTPCGeometry::Row2X(iExtrapolatedRow), tmpY, tmpZ)) { + fail = revert = true; + break; + } + merger.GetConstantMem()->calibObjects.fastTransform->InverseTransformYZtoX(sector, iExtrapolatedRow, tmpY, tmpZ, tmpX); + if (prop.PropagateToXAlpha(tmpX, param.Alpha(sector), inFlyDirection)) { + fail = revert = true; + break; + } + if (CAMath::Abs(mP[2]) > maxSinForUpdate || CAMath::Abs(prop.GetSinPhi0()) >= maxSinForUpdate || mC[0] >= param.rec.tpc.rebuildTrackMaxNonIntCov || mC[2] >= param.rec.tpc.rebuildTrackMaxNonIntCov) { // No further clusters can be attached to the fit + fail = true; + break; + } + if (CAMath::Abs(mP[0]) > CAMath::Abs(mX) * CAMath::Tan(GPUTPCGeometry::kSectAngle() / 2.f) + 0.1f) { + if (iAttempt) { + fail = true; + break; + } + const int32_t sectorSide = sector >= (int8_t)(GPUTPCGeometry::NSECTORS / 2) ? (GPUTPCGeometry::NSECTORS / 2) : 0; + if (mP[0] > 0) { + if (++sector >= sectorSide + 18) { + sector -= 18; } - if GPUCA_RTC_CONSTEXPR (GPUCA_GET_CONSTEXPR(param.rec.tpc, dEdxClusterRejectionFlagMask) != GPUCA_GET_CONSTEXPR(param.rec.tpc, dEdxClusterRejectionFlagMaskAlt)) { - if (accAlt) { - dEdxAlt.fillCluster(qtot, qmax, cluster.row, cluster.sector, mP[2], mP[3], merger->GetConstantMem()->calibObjects, zz, pad, relTime GPUCA_DEBUG_STREAMER_CHECK(, iTrk, clusterState)); - } + } else { + if (--sector < sectorSide) { + sector += 18; } } + if (prop.RotateToAlpha(param.Alpha(sector))) { // GetPropagatedYZ / InverseTransformYZtoX of the next attempt need Y in the frame of the new sector + fail = revert = true; + break; + } + } + } + if (fail) { + break; + } + CADEBUG(printf("\tExtrapol. Sec %2d Row %3d Propaga Alpha %8.3f , X %8.3f - Y %8.3f, Z %8.3f - QPt %7.2f (%7.2f), SP %5.2f (%5.2f) --- Cov sY %8.3f sZ %8.3f sSP %8.3f sPt %8.3f - YPt %8.3f\n", sector, iExtrapolatedRow, prop.GetAlpha(), mX, mP[0], mP[1], mP[4], prop.GetQPt0(), mP[2], prop.GetSinPhi0(), sqrtf(mC[0]), sqrtf(mC[2]), sqrtf(mC[5]), sqrtf(mC[14]), mC[10])); + gputpcgmmergertypes::InterpolationErrorHit inter; + inter.markInvalid(); + float uncorrectedY = FindBestInterpolatedHit(merger, inter, sector, iExtrapolatedRow, deltaZ, sumInvSqrtCharge, nAvgCharge, prop, iTrk, false, true); + auto& candidate = merger.ClusterCandidates()[(iTrk * GPUTPCGeometry::NROWS + iExtrapolatedRow) * param.rec.tpc.rebuildTrackInFitClusterCandidates + 0]; + bool found = candidate.id >= 2; + if (found) { + lastExtrapolateRow = iExtrapolatedRow; +#if defined(DEBUG_REBUILD_MC) && !defined(GPUCA_GPUCODE) + if (merger.MergedTrackMC() && merger.GetConstantMem()->ioPtrs.clustersNative->clustersMCTruth) { + int32_t labelCorrect = GPUTPCTrkLblSearch(merger.GetConstantMem()->ioPtrs.clustersNative->clustersMCTruth->getLabels(candidate.id - 2), merger.MergedTrackMC()[iTrk]); + CADEBUG(printf("\t%21sLabel correct: %d\n", "", labelCorrect)); + } +#endif + float err2Y, err2Z, xx, yy, zz; + const ClusterNative& GPUrestrict() cl = merger.GetConstantMem()->ioPtrs.clustersNative->clustersLinear[candidate.id - 2]; + merger.GetConstantMem()->calibObjects.fastTransform->Transform(sector, iExtrapolatedRow, cl.getPad(), cl.getTime(), xx, yy, zz, mTOffset); + if (prop.PropagateToXAlpha(xx, prop.GetAlpha(), inFlyDirection)) { + candidate.best = -1; + revert = true; + break; + } + const float time = merger.GetConstantMem()->ioPtrs.clustersNative ? cl.getTime() : -1.f; // TODO: When is it possible that we do not have clusterNative? + const float invSqrtCharge = merger.GetConstantMem()->ioPtrs.clustersNative ? CAMath::InvSqrt(cl.qMax) : 0.f; + const float invCharge = merger.GetConstantMem()->ioPtrs.clustersNative ? (1.f / cl.qMax) : 0.f; + float invAvgCharge = (sumInvSqrtCharge += invSqrtCharge) / ++nAvgCharge; + invAvgCharge *= invAvgCharge; + const uint8_t clusterState = cl.getFlags(); + prop.GetErr2(err2Y, err2Z, param, zz, iExtrapolatedRow, clusterState, sector, time, invAvgCharge, invCharge); + CADEBUG(printf("\t%21sResiduals %8.3f %8.3f --- Errors %8.3f %8.3f\n", "", yy - mP[0], zz - mP[1], CAMath::Sqrt(err2Y), CAMath::Sqrt(err2Z))); + if (consecGoodRows == 0) { // New segment + segmentStartChi2 = mChi2; + segmentRows = 0; + } + int32_t retValUpd = prop.Update(yy, zz, iExtrapolatedRow, param, clusterState, true, refit, err2Y, err2Z); + CADEBUG(printf("\tExtrapol. Sec %2d Row %3d Fit Alpha %8.3f , X %8.3f - Y %8.3f, Z %8.3f - QPt %7.2f (%7.2f), SP %5.2f (%5.2f), DzDs %5.2f %16s --- Cov sY %8.3f sZ %8.3f sSP %8.3f sPt %8.3f - YPt %8.3f - FErr %d\n", sector, iExtrapolatedRow, prop.GetAlpha(), mX, mP[0], mP[1], mP[4], prop.GetQPt0(), mP[2], prop.GetSinPhi0(), mP[3], "", sqrtf(mC[0]), sqrtf(mC[2]), sqrtf(mC[5]), sqrtf(mC[14]), mC[10], retValUpd)); + if (retValUpd == 0 && param.rec.tpc.rebuildTrackMaxQPtDrift > 0 && CAMath::Abs(mP[4] - qPtStart) > param.rec.tpc.rebuildTrackMaxQPtDrift * CAMath::Max(1.f, CAMath::Abs(qPtStart))) { // Track changed too much, probably following something else + candidate.best = -1; + revert = true; // The rejected cluster is already in the fit + break; + } + if (retValUpd != 0) { + candidate.best = -1; + if (retValUpd < GPUTPCGMPropagator::updateErrorClusterRejected) { + revert = true; + break; + } + found = false; // Cluster rejected by chi2, handle as missing row + } else { + segmentRows++; } } - } else if (retValInt || retValUpd >= GPUTPCGMPropagator::updateErrorClusterRejected) { // cluster far away form the track - if (retValInt || allowChangeClusters) { - MarkClusters(clusters, ihitMergeFirst, ihit, wayDirection, GPUTPCGMMergedTrackHit::flagRejectDistance); - } else if (finalFit) { - MarkClusters(clusters, ihitMergeFirst, ihit, wayDirection, GPUTPCGMMergedTrackHit::flagRejectErr); + if (found) { + lastFoundRow = iExtrapolatedRow; + lastFoundX = mX; + lastFoundAlpha = prop.GetAlpha(); + if (++consecGoodRows >= param.rec.tpc.rebuildTrackExtrMinConsecGoodRows) { + if (param.rec.tpc.rebuildTrackExtrMaxAvgChi2 <= 0 || mChi2 - segmentStartChi2 <= param.rec.tpc.rebuildTrackExtrMaxAvgChi2 * segmentRows) { + lastGoodRow = iExtrapolatedRow; // TODO: Could optionally store the track state here, and revert to the good state in case we reject further clusters + } + consecGoodRowsMissing = 0; + } + rowGapActive = rowGapTotal = 0; + } else { + if (++missingRowsTotal > param.rec.tpc.rebuildTrackExtrMaxMissingRows) { + break; + } + if (++rowGapTotal > param.rec.tpc.trackFollowingMaxRowGap) { + consecGoodRows = consecGoodRowsMissing = 0; + } + uint32_t pad = CAMath::Float2UIntRn(GPUTPCGeometry::LinearY2Pad(sector, iExtrapolatedRow, uncorrectedY)); + if (pad < GPUTPCGeometry::NPads(iExtrapolatedRow) && (!merger.GetConstantMem()->calibObjects.dEdxCalibContainer || !merger.GetConstantMem()->calibObjects.dEdxCalibContainer->isDead(sector, iExtrapolatedRow, pad))) { + if (rowGapActive++ >= param.rec.tpc.trackFollowingMaxRowGap) { + break; + } + if (consecGoodRowsMissing++ >= param.rec.tpc.rebuildTrackExtrConsecGoodRowsMaxGap) { + consecGoodRows = consecGoodRowsMissing = 0; + } + } } - if (!retValInt) { - nMissed++; - nMissed2++; + } + if (iExtrapolatedRow >= GPUTPCGeometry::NROWS && consecGoodRows >= param.rec.tpc.rebuildTrackExtrMinConsecGoodRowsEdge && (param.rec.tpc.rebuildTrackExtrMaxAvgChi2 <= 0 || mChi2 - segmentStartChi2 <= param.rec.tpc.rebuildTrackExtrMaxAvgChi2 * segmentRows)) { // Reached the inner / outer end of the TPC, a shorter last segment is sufficient + lastGoodRow = lastFoundRow; + } + if (lastGoodRow != lastExtrapolateRow) { + for (int32_t iRow = lastGoodRow + up; iRow != lastExtrapolateRow + up; iRow += up) { + auto& candidate = merger.ClusterCandidates()[(iTrk * GPUTPCGeometry::NROWS + iRow) * param.rec.tpc.rebuildTrackInFitClusterCandidates + 0]; + candidate.best = -1; } - } else { - break; // bad chi2 for the whole track, stop the fit } + if (finalOutInFit && !(param.rec.tpc.disableMarkAdjacent & 4) && iExtrapolatedRow < GPUTPCGeometry::NROWS) { // Continue the looper following from where the extrapolation stopped + StoreLoopPropagation(merger, sector, iExtrapolatedRow, iTrk, up == 1, prop.GetAlpha()); + CADEBUG(printf("\t\tSTORING %d extrapolated row %d sector %d out %d\n", iTrk, (int)iExtrapolatedRow, sector, up == 1)); + } + const int32_t rowsMoved = CAMath::Abs((int32_t)(iExtrapolatedRow < GPUTPCGeometry::NROWS ? iExtrapolatedRow : iExtrapolatedRow - up) - (int32_t)lastGoodRow); + if (revert || lastGoodRow == lastPropagateRow || lastFoundRow != lastGoodRow || (iWay == nWays - 3 && sector != lastSector) || rowsMoved > param.rec.tpc.rebuildTrackExtrMaxRowsBack || (rowsMoved && prop.PropagateToXAlpha(lastFoundX, lastFoundAlpha, !inFlyDirection))) { // Failure, nothing accepted, rejected clusters in the fit, sector changed before the next way goes back, or too far from the last accepted cluster to go back to it + *this = paramStart; + prop.SetTrack(this, alphaStart); + } + } else if (finalOutInFit && !(param.rec.tpc.disableMarkAdjacent & 4) && lastUpdateRow != 255 && !retryAttempt) { + StoreLoopPropagation(merger, lastSector, lastUpdateRow, iTrk, lastUpdateRow > clusters[(iWay & 1) ? (maxN - 1) : 0].row, prop.GetAlpha()); + CADEBUG(printf("\t\tSTORING %d lastUpdateRow %d row %d out %d\n", iTrk, (int)lastUpdateRow, (int)clusters[(iWay & 1) ? (maxN - 1) : 0].row, lastUpdateRow > clusters[(iWay & 1) ? (maxN - 1) : 0].row)); } - if (finalOutInFit && !(merger->Param().rec.tpc.disableRefitAttachment & 4) && lastRow != 255 && lastSector != 255) { - StoreLoopPropagation(merger, lastSector, lastRow, iTrk, lastRow > clusters[(iWay & 1) ? (maxN - 1) : 0].row, prop.GetAlpha()); - CADEBUG(printf("\t\tSTORING %d lastRow %d row %d out %d\n", iTrk, (int)lastRow, (int)clusters[(iWay & 1) ? (maxN - 1) : 0].row, lastRow > clusters[(iWay & 1) ? (maxN - 1) : 0].row)); - } - if (!(iWay & 1) && !finalFit && !track.CCE() && !track.Looper()) { - deltaZ = ShiftZ(clusters, merger, maxN); - } else { - deltaZ = 0.f; + + if (param.rec.tpc.rebuildTrackInFit && iWay == nWays - 2) { + Alpha = prop.GetAlpha(); + if (ihitStart != 0) { + MarkClusters(clusters, 0, ihitStart - 1, 1, GPUTPCGMMergedTrackHit::flagHighIncl); + if (CAMath::Abs(clusters[ihitStart].row - clusters[0].row) > param.rec.tpc.rebuildTrackMaxRowsFromEnd) { // The fit ended far away from the end of the track, the final fit cannot start from there, so let it start from the original parameters like the first way + const int32_t ndf = mNDF; + const float chi2 = mChi2; + *this = track.GetParam(); + mNDF = ndf; + mChi2 = chi2; + Alpha = track.GetAlpha(); + } + } + return true; } } ConstrainSinPhi(); @@ -344,7 +496,7 @@ GPUd() bool GPUTPCGMTrackParam::Fit(GPUTPCGMMerger* GPUrestrict() merger, int32_ o2::utils::DebugStreamer::instance()->getStreamer("debug_accept_track", "UPDATE") << o2::utils::DebugStreamer::instance()->getUniqueTreeName("debug_accept_track").data() << "iTrk=" << iTrk << "outerParam=" << track.OuterParam() << "track=" << this << "ihitStart=" << ihitStart << "\n"; }) - if (!(N + NTolerated >= (int32_t)param.tpcMinHitsB5(mP[4] * param.qptB5Scaler) && 2 * NTolerated <= CAMath::Max(10, N) && CheckNumericalQuality(covYYUpd))) { + if (!(N + NTolerated >= (int32_t)merger.MinHitsForQPt(mP[4], mC[14]) && 2 * NTolerated <= CAMath::Max(10, N) && CheckNumericalQuality(covYYUpd))) { return false; // TODO: NTolerated should never become that large, check what is going wrong! } if (param.rec.tpc.minNClustersFinalTrack != -1 && N + NTolerated < param.rec.tpc.minNClustersFinalTrack) { @@ -352,9 +504,9 @@ GPUd() bool GPUTPCGMTrackParam::Fit(GPUTPCGMMerger* GPUrestrict() merger, int32_ } if (param.par.dodEdx && param.dodEdxEnabled) { - dEdx.computedEdx(merger->MergedTracksdEdx()[iTrk], param); + dEdx.computedEdx(merger.MergedTracksdEdx()[iTrk], param); if GPUCA_RTC_CONSTEXPR (GPUCA_GET_CONSTEXPR(param.rec.tpc, dEdxClusterRejectionFlagMask) != GPUCA_GET_CONSTEXPR(param.rec.tpc, dEdxClusterRejectionFlagMaskAlt)) { - dEdxAlt.computedEdx(merger->MergedTracksdEdxAlt()[iTrk], param); + dEdxAlt.computedEdx(merger.MergedTracksdEdxAlt()[iTrk], param); } } Alpha = prop.GetAlpha(); @@ -364,37 +516,371 @@ GPUd() bool GPUTPCGMTrackParam::Fit(GPUTPCGMMerger* GPUrestrict() merger, int32_ return true; } -GPUdni() void GPUTPCGMTrackParam::MoveToReference(GPUTPCGMPropagator& prop, const GPUParam& param, float& Alpha) +GPUdii() void GPUTPCGMTrackParam::FitAddRow(const int32_t iRow, const uint8_t sector, const int32_t iTrk, const GPUTPCGMMergedTrack& GPUrestrict() track, GPUTPCGMPropagator& GPUrestrict() prop, const bool inFlyDirection, GPUTPCGMMerger& GPUrestrict() merger, uint8_t* GPUrestrict() dEdxSubThresholdRow, const bool dodEdx, const bool doAttach, const uint8_t doInterpolate, gputpcgmmergertypes::InterpolationErrorHit& GPUrestrict() inter, const float deltaZ, const float sumInvSqrtCharge, const int32_t nAvgCharge) +{ + if (CAMath::Abs(mP[2]) > constants::MAX_SIN_PHI || CAMath::Abs(mP[0]) > CAMath::Abs(mX) * CAMath::Tan(GPUTPCGeometry::kSectAngle() / 2.f)) { + return; + } + const GPUParam& GPUrestrict() param = merger.Param(); + if GPUCA_RTC_CONSTEXPR (GPUCA_GET_CONSTEXPR(param.par, dodEdx)) { + if (dodEdx) { + float yUncorrected, zUncorrected; + merger.GetConstantMem()->calibObjects.fastTransform->InverseTransformYZtoNominalYZ(sector, iRow, mP[0], mP[1], yUncorrected, zUncorrected); + uint32_t pad = CAMath::Float2UIntRn(GPUTPCGeometry::LinearY2Pad(sector, iRow, yUncorrected)); + if (!(pad >= GPUTPCGeometry::NPads(iRow) || (merger.GetConstantMem()->calibObjects.dEdxCalibContainer && merger.GetConstantMem()->calibObjects.dEdxCalibContainer->isDead(sector, iRow, pad)))) { + *dEdxSubThresholdRow = iRow; + } + } + } + if (doAttach) { + AttachClusters(merger, sector, iRow, iTrk, track.Leg() == 0, prop); + } + if (doInterpolate) { + if (doInterpolate == 2) { + if (inter.isValid() && CAMath::Abs(inter.posY - mP[0]) > param.rec.tpc.rebuildTrackMaxInterpolationDistY) { + inter.markInvalid(); + } + FindBestInterpolatedHit(merger, inter, sector, iRow, deltaZ, sumInvSqrtCharge, nAvgCharge, prop, iTrk, false); + } else { + prop.InterpolateFill(&inter); + } + } +} + +GPUdii() void GPUTPCGMTrackParam::HandleCrossCE(const GPUParam& GPUrestrict() param, const uint8_t sector, const uint8_t& lastSector) { - constexpr float kDeg2Rad = M_PI / 180.f; - constexpr float kSectAngle = 2 * M_PI / 18.f; + const bool crossCE = lastSector != 255 && ((lastSector < 18) ^ (sector < 18)); + if (crossCE) { + if (param.rec.tpc.addErrorsCECrossing) { + if (param.rec.tpc.addErrorsCECrossing >= 2) { + AddCovDiagErrorsWithCorrelations(param.rec.tpc.errorsCECrossing); + } else { + AddCovDiagErrors(param.rec.tpc.errorsCECrossing); + } + } else if (mC[2] < 0.5f) { + mC[2] = 0.5f; + } + } +} +GPUdii() int32_t GPUTPCGMTrackParam::FitHit(GPUTPCGMMerger& GPUrestrict() merger, const int32_t iTrk, const GPUTPCGMMergedTrack& GPUrestrict() track, const float xx, const float yy, const float zz, const uint8_t clusterState, const float clAlpha, const int32_t iWay, const bool inFlyDirection, float& GPUrestrict() deltaZ, float& GPUrestrict() lastUpdateX, GPUTPCGMMergedTrackHit* GPUrestrict() clusters, GPUTPCGMPropagator& GPUrestrict() prop, gputpcgmmergertypes::InterpolationErrorHit& GPUrestrict() inter, GPUdEdx& GPUrestrict() dEdx, GPUdEdx& GPUrestrict() dEdxAlt, float& GPUrestrict() sumInvSqrtCharge, int32_t& GPUrestrict() nAvgCharge, const int32_t ihit, const int32_t ihitMergeFirst, const bool allowChangeClusters, const bool refit, const bool finalFit, int32_t& GPUrestrict() nMissed, int32_t& GPUrestrict() nMissed2, float uncorrectedY, bool retryAttempt) +{ + const GPUParam& GPUrestrict() param = merger.Param(); + const int32_t nWays = param.rec.tpc.nWays; + const int32_t wayDirection = (iWay & 1) ? -1 : 1; + const auto& cluster = clusters[ihit]; + + int32_t retValUpd = 0, retValInt = 0; + float threshold = 3.f + (lastUpdateX >= 0 ? (CAMath::Abs(mX - lastUpdateX) / 2) : 0.f); + if (mNDF > (int32_t)param.rec.tpc.mergerNonInterpolateRejectMinNDF && (CAMath::Abs(yy - mP[0]) > threshold || CAMath::Abs(zz - mP[1]) > threshold)) { + retValUpd = GPUTPCGMPropagator::updateErrorClusterRejectedDistance; + if (param.rec.tpc.rebuildTrackInFit) { + MarkClusters(clusters, ihitMergeFirst, ihit, wayDirection, GPUTPCGMMergedTrackHit::flagRejectDistance); // Will enable rejectChi2 in further rounds + } + } else { + float err2Y, err2Z; + const float time = merger.GetConstantMem()->ioPtrs.clustersNative ? merger.GetConstantMem()->ioPtrs.clustersNative->clustersLinear[cluster.num].getTime() : -1.f; + const float invSqrtCharge = merger.GetConstantMem()->ioPtrs.clustersNative ? CAMath::InvSqrt(merger.GetConstantMem()->ioPtrs.clustersNative->clustersLinear[cluster.num].qMax) : 0.f; + const float invCharge = merger.GetConstantMem()->ioPtrs.clustersNative ? (1.f / merger.GetConstantMem()->ioPtrs.clustersNative->clustersLinear[cluster.num].qMax) : 0.f; + float invAvgCharge = (sumInvSqrtCharge += invSqrtCharge) / ++nAvgCharge; + invAvgCharge *= invAvgCharge; + + prop.GetErr2(err2Y, err2Z, param, zz, cluster.row, clusterState, cluster.sector, time, invAvgCharge, invCharge); + + bool rejectChi2 = !retryAttempt && (clusterState & GPUTPCGMMergedTrackHit::flagReject); + if (!retryAttempt) { + if (param.rec.tpc.mergerInterpolateErrors) { + if (iWay == nWays - 2) { + if (!param.rec.tpc.rebuildTrackInFit) { + if (inter.isValid()) { + retValInt = prop.InterpolateReject(param, yy, zz, clusterState, &inter, err2Y, err2Z, deltaZ); + } else { + rejectChi2 = true; + } + } else if (track.NClusters() <= param.rec.tpc.rebuildTrackFallbackMaxNCl && inter.isValid()) { // Short track: keep the original cluster as candidate if the search found none, but it passes the interpolated chi2 cut + auto& candidate = merger.ClusterCandidates()[(iTrk * GPUTPCGeometry::NROWS + cluster.row) * param.rec.tpc.rebuildTrackInFitClusterCandidates]; + if (candidate.id < 2 && !prop.InterpolateReject(param, yy, zz, clusterState, &inter, err2Y, err2Z, deltaZ)) { + const float wY = mC[0] / (mC[0] + (float)inter.errorY), wZ = mC[2] / (mC[2] + (float)inter.errorZ); // Weight of the stored position in the interpolation + candidate = {.id = cluster.num + 2, .row = cluster.row, .sector = cluster.sector, .error = CAMath::Square(yy - mP[0] - wY * ((float)inter.posY - mP[0])) + CAMath::Square(zz - mP[1] - wZ * ((float)inter.posZ + deltaZ - mP[1])), .weight = 0, .best = 0}; + } + } + } else if (iWay == nWays - 1) { + if (param.rec.tpc.mergerInterpolateRejectAlsoOnCurrentPosition && GetNDF() > (int32_t)param.rec.tpc.mergerNonInterpolateRejectMinNDF) { + rejectChi2 = true; + } + } + } else { + rejectChi2 = allowChangeClusters; + } + } + + if (param.rec.tpc.rejectEdgeClustersInTrackFit && uncorrectedY > -1e6f && param.rejectEdgeClusterByY(uncorrectedY, cluster.row, CAMath::Sqrt(mC[0]))) { + retValUpd = GPUTPCGMPropagator::updateErrorClusterRejectedEdge; + } else { + retValUpd = prop.Update(yy, zz, cluster.row, param, clusterState, rejectChi2, refit, err2Y, err2Z); + } + GPUCA_DEBUG_STREAMER_CHECK(if (o2::utils::DebugStreamer::checkStream(o2::utils::StreamFlags::streamUpdateTrack, iTrk)) { + merger.DebugStreamerUpdate(iTrk, ihit, xx, yy, zz, cluster, merger.GetConstantMem()->ioPtrs.clustersNative->clustersLinear[cluster.num], *this, prop, inter, rejectChi2, refit, retValUpd, sumInvSqrtCharge / nAvgCharge * sumInvSqrtCharge / nAvgCharge, yy, zz, clusterState, retValInt, err2Y, err2Z); + }); + } + // clang-format off + CADEBUG(if (!CheckCov()) GPUError("INVALID COV AFTER UPDATE!!!")); + CADEBUG(printf("\t%21sFit Alpha %8.3f , X %8.3f - Y %8.3f, Z %8.3f - QPt %7.2f (%7.2f), SP %5.2f (%5.2f), DzDs %5.2f %16s --- Cov sY %8.3f sZ %8.3f sSP %8.3f sPt %8.3f - YPt %8.3f - FErr %d %d\n", "", prop.GetAlpha(), mX, mP[0], mP[1], mP[4], prop.GetQPt0(), mP[2], prop.GetSinPhi0(), mP[3], "", sqrtf(mC[0]), sqrtf(mC[2]), sqrtf(mC[5]), sqrtf(mC[14]), mC[10], retValUpd, retValInt)); + // clang-format on + + ConstrainSinPhi(); // TODO: Limit using ConstrainSinPhi everywhere! + if (!retValUpd && !retValInt) // track is updated + { + lastUpdateX = mX; + nMissed = nMissed2 = 0; + UnmarkClusters(clusters, ihitMergeFirst, ihit, wayDirection, GPUTPCGMMergedTrackHit::flagHighIncl); + return 0; // ok + } else if (retValInt || retValUpd >= GPUTPCGMPropagator::updateErrorClusterRejected) { // cluster far away form the track + if (retValInt || allowChangeClusters) { + MarkClusters(clusters, ihitMergeFirst, ihit, wayDirection, GPUTPCGMMergedTrackHit::flagRejectDistance); + } else if (finalFit) { + MarkClusters(clusters, ihitMergeFirst, ihit, wayDirection, retValUpd >= GPUTPCGMPropagator::updateErrorClusterRejected ? GPUTPCGMMergedTrackHit::flagRejectDistance : GPUTPCGMMergedTrackHit::flagRejectErr); + } + if (!retValInt) { + nMissed++; + nMissed2++; + } + return -1; // cluster rejected + } else { + return 1; // bad chi2 for the whole track, stop the fit + } +} + +GPUdii() float GPUTPCGMTrackParam::FindBestInterpolatedHit(GPUTPCGMMerger& GPUrestrict() merger, gputpcgmmergertypes::InterpolationErrorHit& GPUrestrict() inter, const uint8_t sector, const uint8_t row, const float deltaZ, const float sumInvSqrtCharge, const int nAvgCharge, const GPUTPCGMPropagator& GPUrestrict() prop, const int32_t iTrk, bool interOnly, bool extrapolation) +{ + const GPUParam& GPUrestrict() param = merger.Param(); + const GPUTPCTracker& GPUrestrict() tracker = *(merger.GetConstantMem()->tpcTrackers + sector); + const GPUTPCRow& GPUrestrict() rowData = tracker.Row(row); + GPUglobalref() const cahit2* hits = tracker.HitData(rowData); + GPUglobalref() const calink* firsthit = tracker.FirstHitInBin(rowData); + float uncorrectedY = -1e6f, uncorrectedZ; + const float maxNonIntCov = param.rec.tpc.rebuildTrackMaxNonIntCov; + bool useInter = inter.isValid(); + if (useInter && interOnly) { // One-sided from the other direction, same limit as for the current track position + useInter = (float)inter.errorY < maxNonIntCov && (float)inter.errorZ < maxNonIntCov; + } + const bool useCurrent = !interOnly && maxNonIntCov > 0 && mC[0] < maxNonIntCov && mC[2] < maxNonIntCov && mNDF > (int32_t)param.rec.tpc.rebuildTrackNonIntMinNDF; + if (rowData.NHits() && (useInter || useCurrent)) { + const float zOffset = param.par.continuousTracking ? merger.GetConstantMem()->calibObjects.fastTransform->convVertexTimeToZOffset(sector, mTOffset, param.continuousMaxTimeBin) : 0; + const float y0 = rowData.Grid().YMin(); + const float stepY = rowData.HstepY(); + const float z0 = rowData.Grid().ZMin() - zOffset; // We can use our own ZOffset, since this is only used temporarily anyway + const float stepZ = rowData.HstepZ(); + int32_t bin, ny, nz; + + float ImP0, ImP1, ImC0, ImC2; + if (interOnly) { + ImP0 = (float)inter.posY; + ImP1 = (float)inter.posZ + deltaZ; + ImC0 = (float)inter.errorY; + ImC2 = (float)inter.errorZ; + } else if (useInter) { + const float Iz0 = (float)inter.posY - mP[0]; + const float Iz1 = (float)inter.posZ + deltaZ - mP[1]; + const float Iw0 = 1.f / (mC[0] + (float)inter.errorY); + const float Iw2 = 1.f / (mC[2] + (float)inter.errorZ); + const float Ik00 = mC[0] * Iw0; + const float Ik11 = mC[2] * Iw2; + ImP0 = mP[0] + Ik00 * Iz0; + ImP1 = mP[1] + Ik11 * Iz1; + ImC0 = mC[0] - Ik00 * mC[0]; + ImC2 = mC[2] - Ik11 * mC[2]; + } else { + ImP0 = mP[0]; + ImP1 = mP[1]; + ImC0 = mC[0]; + ImC2 = mC[2]; + } + + merger.GetConstantMem()->calibObjects.fastTransform->InverseTransformYZtoNominalYZ(sector, row, ImP0, ImP1, uncorrectedY, uncorrectedZ); + + float xMoved = mX, yMoved = mP[0], zMoved = mP[1], sinPhiMoved = prop.GetSinPhi0(); + if (interOnly) { // Track is not at this row: estimate track position and angle there (for the cluster errors) from the interpolated position and the helix of the current linearization + float sA, cA; + CAMath::SinCos(param.Alpha(sector) - prop.GetAlpha(), sA, cA); + xMoved = GPUTPCGeometry::Row2X(row); + yMoved = ImP0; + zMoved = ImP1; + const float sinPhi = prop.GetSinPhi0() - prop.GetQPt0() * prop.GetBz(mX, mP[0], mP[1]) * (xMoved * cA - yMoved * sA - mX); // at the X of the row in the current track frame + if (CAMath::Abs(sinPhi) < constants::MAX_SIN_PHI) { + sinPhiMoved = sinPhi * cA - CAMath::Copysign(CAMath::Sqrt(1.f - sinPhi * sinPhi), prop.GetCosPhi0()) * sA; + } + sinPhiMoved = CAMath::Clamp(sinPhiMoved, -constants::MAX_SIN_PHI, constants::MAX_SIN_PHI); + } + + int32_t nCandidates = 0; + while (nCandidates < param.rec.tpc.rebuildTrackInFitClusterCandidates && merger.ClusterCandidates()[(iTrk * GPUTPCGeometry::NROWS + row) * param.rec.tpc.rebuildTrackInFitClusterCandidates + nCandidates].id > 1) { + nCandidates++; + } + if (CAMath::Abs(uncorrectedY) <= rowData.getTPCMaxY()) { + float err2Y, err2Z; + param.GetClusterErrors2(sector, row, zMoved, sinPhiMoved, mP[3], -1.f, 0.f, 0.f, err2Y, err2Z); // TODO: Use correct time/avgCharge + + const float kFactor = tracker.GetChiSeedFactor(); + const float sy2 = 4 * CAMath::Min(param.rec.tpc.hitSearchArea2, kFactor * (err2Y + CAMath::Abs(ImC0))); // TODO: is 4 a good factor?? + const float sz2 = 4 * CAMath::Min(param.rec.tpc.hitSearchArea2, kFactor * (err2Z + CAMath::Abs(ImC2))); + const float tubeY = CAMath::Sqrt(sy2); + const float tubeZ = CAMath::Sqrt(sz2); + rowData.Grid().GetBinArea(uncorrectedY, uncorrectedZ + zOffset, tubeY, tubeZ, bin, ny, nz); + + const float errFactor = extrapolation ? param.rec.tpc.rebuildTrackErrFactorExtrapolation : ((interOnly || !useInter) ? param.rec.tpc.rebuildTrackErrFactorOneSided : 1.f); // Downweight candidates not found by two-sided interpolation + const int32_t nBinsY = rowData.Grid().Ny(); + const int32_t idOffset = tracker.Data().ClusterIdOffset(); + const int32_t* ids = &(tracker.Data().ClusterDataIndex()[rowData.HitNumberOffset()]); + for (int32_t k = 0; k <= nz; k++) { + const int32_t mybin = bin + k * nBinsY; + const uint32_t hitFst = firsthit[mybin]; + const uint32_t hitLst = firsthit[mybin + ny + 1]; + for (uint32_t ih = hitFst; ih < hitLst; ih++) { + const cahit2 hh = hits[ih]; + const float y = y0 + hh.x * stepY; + const float z = z0 + hh.y * stepZ; + const float dy = y - uncorrectedY; + const float dz = z - uncorrectedZ; + + if (dy * dy < sy2 && dz * dz < sz2) { + float err2YA, err2ZA; + const ClusterNative& GPUrestrict() cl = merger.GetConstantMem()->ioPtrs.clustersNative->clustersLinear[idOffset + ids[ih]]; + const auto clflags = cl.getFlags() & GPUTPCGMMergedTrackHit::clustererAndSharedFlags; + const float time = cl.getTime(); + const float invSqrtCharge = CAMath::InvSqrt(cl.qMax); + const float invCharge = 1.f / cl.qMax; + float invAvgCharge = (sumInvSqrtCharge + invSqrtCharge) / (nAvgCharge + 1); + invAvgCharge *= invAvgCharge; + + if (interOnly) { + GPUTPCGMPropagator::GetErr2(err2YA, err2ZA, param, sinPhiMoved, mP[3], zMoved, xMoved, yMoved, row, clflags, sector, time, invAvgCharge, invCharge, false); + } else { + prop.GetErr2(err2YA, err2ZA, param, mP[1], row, clflags, sector, time, invAvgCharge, invCharge); + } + const float Jw0 = 1.f / (ImC0 + err2YA); + const float Jw2 = 1.f / (ImC2 + err2ZA); + const float chi2Y = Jw0 * dy * dy; + const float chi2Z = Jw2 * dz * dz; + bool ok = !prop.RejectCluster(chi2Y * param.rec.tpc.clusterRejectChi2TolleranceY, chi2Z * param.rec.tpc.clusterRejectChi2TolleranceZ, clflags); + const float err = (param.rec.tpc.rebuildTrackCandidateErrChi2 ? (chi2Y + chi2Z) : (dy * dy + dz * dz)) * errFactor; + if (ok) { + int32_t insert = nCandidates; + for (int32_t c = 0; c < nCandidates; c++) { + if (err < merger.ClusterCandidates()[(iTrk * GPUTPCGeometry::NROWS + row) * param.rec.tpc.rebuildTrackInFitClusterCandidates + c].error) { + insert = c; + break; + } + } + if (insert < param.rec.tpc.rebuildTrackInFitClusterCandidates) { + for (int32_t c = CAMath::Min(nCandidates, param.rec.tpc.rebuildTrackInFitClusterCandidates - 1); c > insert; c--) { + merger.ClusterCandidates()[(iTrk * GPUTPCGeometry::NROWS + row) * param.rec.tpc.rebuildTrackInFitClusterCandidates + c] = merger.ClusterCandidates()[(iTrk * GPUTPCGeometry::NROWS + row) * param.rec.tpc.rebuildTrackInFitClusterCandidates + c - 1]; + } // TODO: Downweight non-interpolated candidates + merger.ClusterCandidates()[(iTrk * GPUTPCGeometry::NROWS + row) * param.rec.tpc.rebuildTrackInFitClusterCandidates + insert] = {.id = (uint32_t)(idOffset + ids[ih] + 2), .row = row, .sector = sector, .error = err, .weight = 0, .best = 0}; + nCandidates += (nCandidates < param.rec.tpc.rebuildTrackInFitClusterCandidates); + } + } + } + } + } + CADEBUG(const auto* dbgCand = &merger.ClusterCandidates()[(iTrk * GPUTPCGeometry::NROWS + row) * param.rec.tpc.rebuildTrackInFitClusterCandidates]; for (int dbgi = 0; dbgi < nCandidates; dbgi++) { if (dbgCand[dbgi].id > 1) printf("\t\t\tiTrk %d Row %d Candidate %d hit %d err %f\n", iTrk, (int)row, dbgi, dbgCand[dbgi].id - 2, dbgCand[dbgi].error); else break; }); + } + if (nCandidates == 0 && !interOnly) { // TODO: remove the interOnly here when rebuilding is fully working + merger.ClusterCandidates()[(iTrk * GPUTPCGeometry::NROWS + row) * param.rec.tpc.rebuildTrackInFitClusterCandidates + 0].id = 1; + } + } + return uncorrectedY; +} + +GPUdni() void GPUTPCGMTrackParam::InterpolateMissingRows(GPUTPCGMMerger& GPUrestrict() merger, gputpcgmmergertypes::InterpolationErrors& GPUrestrict() interpolation, GPUTPCGMMergedTrackHit* GPUrestrict() clusters, int32_t ihit, int32_t interpolationIndex, int32_t lastRow, const float deltaZ, const float sumInvSqrtCharge, const int32_t nAvgCharge, const GPUTPCGMPropagator& GPUrestrict() prop, const int32_t iTrk) +{ + for (; ihit >= 0; ihit--, interpolationIndex--) { + while (ihit > 0 && clusters[ihit].row == clusters[ihit - 1].row && clusters[ihit].sector == clusters[ihit - 1].sector) { + ihit--; + } + const auto& cluster = clusters[ihit]; + if (CAMath::Abs(cluster.row - lastRow) > 1) { + interpolationIndex -= CAMath::Abs(cluster.row - lastRow) - 1; + } + lastRow = cluster.row; + if (interpolationIndex < 0) { // TODO: Is this check needed? + return; + } + auto inter = interpolation.hit[interpolationIndex]; + FindBestInterpolatedHit(merger, inter, cluster.sector, cluster.row, deltaZ, sumInvSqrtCharge, nAvgCharge, prop, iTrk, true); + } +} + +GPUd() void GPUTPCGMTrackParam::DodEdx(GPUdEdx& GPUrestrict() dEdx, GPUdEdx& GPUrestrict() dEdxAlt, GPUTPCGMMerger& GPUrestrict() merger, bool finalFit, int ihit, int ihitMergeFirst, int wayDirection, const GPUTPCGMMergedTrackHit* GPUrestrict() clusters, uint8_t clusterState, float zz, uint8_t dEdxSubThresholdRow GPUCA_DEBUG_STREAMER_CHECK(, int32_t iTrk)) +{ + const GPUParam& GPUrestrict() param = merger.Param(); + if GPUCA_RTC_CONSTEXPR (GPUCA_GET_CONSTEXPR(param.par, dodEdx)) { + const GPUCalibObjectsConst& GPUrestrict() calib = merger.GetConstantMem()->calibObjects; + const ClusterNative* GPUrestrict() clustersArray = merger.GetConstantMem()->ioPtrs.clustersNative->clustersLinear; + if (param.dodEdxEnabled && finalFit) { // TODO: Costimize flag to remove, and option to remove double-clusters + bool acc = (clusterState & param.rec.tpc.dEdxClusterRejectionFlagMask) == 0, accAlt = (clusterState & param.rec.tpc.dEdxClusterRejectionFlagMaskAlt) == 0; + if (acc || accAlt) { + float qtot = 0, qmax = 0, pad = 0, relTime = 0; + const int32_t clusterCount = CAMath::Abs(ihit - ihitMergeFirst) + 1; + for (int32_t iTmp = ihitMergeFirst; iTmp != ihit + wayDirection; iTmp += wayDirection) { + const ClusterNative& cl = clustersArray[clusters[iTmp].num]; + qtot += cl.getQtot(); + qmax = CAMath::Max(qmax, cl.qMax); + pad += cl.getPad(); + relTime += cl.getTime(); + } + qtot /= clusterCount; // TODO: Weighted Average + pad /= clusterCount; + relTime /= clusterCount; + relTime = relTime - CAMath::Round(relTime); + const auto& cluster = clusters[ihit]; + if (acc) { + if (dEdxSubThresholdRow != 255) { // Must be called before fillCluster, which flushes pending sub-threshold rows with the min charge of their ROC when the ROC changes + dEdx.fillSubThreshold(dEdxSubThresholdRow); + } + dEdx.fillCluster(qtot, qmax, cluster.row, cluster.sector, mP[2], mP[3], calib, zz, pad, relTime GPUCA_DEBUG_STREAMER_CHECK(, iTrk, clusterState)); + } + if GPUCA_RTC_CONSTEXPR (GPUCA_GET_CONSTEXPR(param.rec.tpc, dEdxClusterRejectionFlagMask) != GPUCA_GET_CONSTEXPR(param.rec.tpc, dEdxClusterRejectionFlagMaskAlt)) { + if (accAlt) { + if (dEdxSubThresholdRow != 255) { + dEdxAlt.fillSubThreshold(dEdxSubThresholdRow); + } + dEdxAlt.fillCluster(qtot, qmax, cluster.row, cluster.sector, mP[2], mP[3], calib, zz, pad, relTime GPUCA_DEBUG_STREAMER_CHECK(, iTrk, clusterState)); + } + } + } + } + } +} + +GPUdni() void GPUTPCGMTrackParam::MoveToReference(GPUTPCGMPropagator& GPUrestrict() prop, const GPUParam& GPUrestrict() param, float& GPUrestrict() Alpha) +{ if (param.rec.tpc.trackReferenceX <= 500) { GPUTPCGMTrackParam save = *this; float saveAlpha = Alpha; for (int32_t attempt = 0; attempt < 3; attempt++) { - float dAngle = CAMath::Round(CAMath::ATan2(mP[0], mX) / kDeg2Rad / 20.f) * kSectAngle; + float dAngle = CAMath::Round(CAMath::ATan2(mP[0], mX) / CAMath::Deg2Rad() / 20.f) * GPUTPCGeometry::kSectAngle(); Alpha += dAngle; if (prop.PropagateToXAlpha(param.rec.tpc.trackReferenceX, Alpha, 0)) { break; } ConstrainSinPhi(); - if (CAMath::Abs(mP[0]) <= mX * CAMath::Tan(kSectAngle / 2.f)) { + if (CAMath::Abs(mP[0]) <= mX * CAMath::Tan(GPUTPCGeometry::kSectAngle() / 2.f)) { return; } } *this = save; Alpha = saveAlpha; } - if (CAMath::Abs(mP[0]) > mX * CAMath::Tan(kSectAngle / 2.f)) { - float dAngle = CAMath::Round(CAMath::ATan2(mP[0], mX) / kDeg2Rad / 20.f) * kSectAngle; + if (CAMath::Abs(mP[0]) > mX * CAMath::Tan(GPUTPCGeometry::kSectAngle() / 2.f)) { + float dAngle = CAMath::Round(CAMath::ATan2(mP[0], mX) / CAMath::Deg2Rad() / 20.f) * GPUTPCGeometry::kSectAngle(); Rotate(dAngle); ConstrainSinPhi(); Alpha += dAngle; } } -GPUd() void GPUTPCGMTrackParam::MirrorTo(GPUTPCGMPropagator& GPUrestrict() prop, float toY, float toZ, bool inFlyDirection, const GPUParam& param, uint8_t row, uint8_t clusterState, bool mirrorParameters, int8_t sector) +GPUd() void GPUTPCGMTrackParam::MirrorTo(GPUTPCGMPropagator& GPUrestrict() prop, float toY, float toZ, bool inFlyDirection, const GPUParam& GPUrestrict() param, uint8_t row, uint8_t clusterState, bool mirrorParameters, int8_t sector) { if (mirrorParameters) { prop.Mirror(inFlyDirection); @@ -421,79 +907,110 @@ GPUd() void GPUTPCGMTrackParam::MirrorTo(GPUTPCGMPropagator& GPUrestrict() prop, mChi2 = 0; } -GPUd() int32_t GPUTPCGMTrackParam::MergeDoubleRowClusters(int32_t& ihit, int32_t wayDirection, GPUTPCGMMergedTrackHit* GPUrestrict() clusters, const GPUTPCGMMerger* GPUrestrict() merger, GPUTPCGMPropagator& GPUrestrict() prop, float& GPUrestrict() xx, float& GPUrestrict() yy, float& GPUrestrict() zz, int32_t maxN, float clAlpha, uint8_t& GPUrestrict() clusterState, bool rejectChi2) +GPUd() int32_t GPUTPCGMTrackParam::MergeDoubleRowClusters(int32_t& ihit, int32_t wayDirection, GPUTPCGMMergedTrackHit* GPUrestrict() clusters, const GPUTPCGMMerger& GPUrestrict() merger, GPUTPCGMPropagator& GPUrestrict() prop, float& GPUrestrict() xx, float& GPUrestrict() yy, float& GPUrestrict() zz, int32_t maxN, float clAlpha, uint8_t& GPUrestrict() clusterState, const bool markReject) { + const int32_t ihitFirst = ihit; + { + const ClusterNative& GPUrestrict() cl = merger.GetConstantMem()->ioPtrs.clustersNative->clustersLinear[clusters[ihit].num]; + merger.GetConstantMem()->calibObjects.fastTransform->Transform(clusters[ihit].sector, clusters[ihit].row, cl.getPad(), cl.getTime(), xx, yy, zz, mTOffset); + } if (ihit + wayDirection >= 0 && ihit + wayDirection < maxN && clusters[ihit].row == clusters[ihit + wayDirection].row && clusters[ihit].sector == clusters[ihit + wayDirection].sector) { - float maxDistY, maxDistZ; - prop.GetErr2(maxDistY, maxDistZ, merger->Param(), zz, clusters[ihit].row, 0, clusters[ihit].sector, -1.f, 0.f, 0.f); // TODO: Use correct time, avgCharge - maxDistY = (maxDistY + mC[0]) * 20.f; - maxDistZ = (maxDistZ + mC[2]) * 20.f; - int32_t noReject = 0; // Cannot reject if simple estimation of y/z fails (extremely unlike case) + float maxDistY2, maxDistZ2; + bool noReject = false; // Cannot reject if simple estimation of y/z fails (extremely unlike case) if (CAMath::Abs(clAlpha - prop.GetAlpha()) > 1.e-4f) { noReject = prop.RotateToAlpha(clAlpha); } float projY = 0, projZ = 0; - if (noReject == 0) { + if (!noReject) { noReject |= prop.GetPropagatedYZ(xx, projY, projZ); } - float count = 0.f; - xx = yy = zz = 0.f; - clusterState = 0; - while (true) { - const ClusterNative& GPUrestrict() cl = merger->GetConstantMem()->ioPtrs.clustersNative->clustersLinear[clusters[ihit].num]; - float clamp = cl.getQtot(); - float clx, cly, clz; - merger->GetConstantMem()->calibObjects.fastTransform->Transform(clusters[ihit].sector, clusters[ihit].row, cl.getPad(), cl.getTime(), clx, cly, clz, mTOffset); - float dy = cly - projY; - float dz = clz - projZ; - if (noReject == 0 && (dy * dy > maxDistY || dz * dz > maxDistZ)) { - CADEBUG(printf("\t\tRejecting double-row cluster: dy %f, dz %f, chiY %f, chiZ %f (Y: trk %f prj %f cl %f - Z: trk %f prj %f cl %f)\n", dy, dz, sqrtf(maxDistY), sqrtf(maxDistZ), mP[0], projY, cly, mP[1], projZ, clz)); - if (rejectChi2) { - clusters[ihit].state |= GPUTPCGMMergedTrackHit::flagRejectDistance; + prop.GetErr2(maxDistY2, maxDistZ2, merger.Param(), zz, clusters[ihit].row, 0, clusters[ihit].sector, -1.f, 0.f, 0.f); // TODO: Use correct time, avgCharge + const float kFactor = merger.GetConstantMem()->tpcTrackers[0].GetChiSeedFactor() * 4.f; + maxDistY2 = (maxDistY2 + mC[0]) * kFactor; + maxDistZ2 = (maxDistZ2 + mC[2]) * kFactor; + auto chkFunction = [clusters, markReject, maxDistY2, maxDistZ2, projY, projZ, noReject CADEBUG3(, this)](int32_t ih, float y, float z) { + float dy = y - projY; + float dz = z - projZ; + if (!noReject && (dy * dy > maxDistY2 || dz * dz > maxDistZ2)) { + CADEBUG(printf("\t\t\tRejecting double-row cluster: dy %f, dz %f, chiY %f, chiZ %f (Y: trk %f prj %f cl %f - Z: trk %f prj %f cl %f)\n", dy, dz, sqrtf(maxDistY2), sqrtf(maxDistZ2), mP[0], projY, y, mP[1], projZ, z)); + if (markReject) { + clusters[ih].state |= GPUTPCGMMergedTrackHit::flagRejectDistance; } + return false; } else { - CADEBUG(printf("\t\tMerging hit row %d X %f Y %f Z %f (dy %f, dz %f, chiY %f, chiZ %f)\n", clusters[ihit].row, clx, cly, clz, dy, dz, sqrtf(maxDistY), sqrtf(maxDistZ))); - xx += clx * clamp; // TODO: Weight in pad/time instead of XYZ + CADEBUG(printf("\t\t\tMerging hit row %d Y %f Z %f (dy %f, dz %f, chiY %f, chiZ %f)\n", clusters[ih].row, y, z, dy, dz, sqrtf(maxDistY2), sqrtf(maxDistZ2))); + return true; + } + }; + const float tmpX = xx; + float count; + if (chkFunction(ihit, yy, zz)) { + const ClusterNative& GPUrestrict() cl = merger.GetConstantMem()->ioPtrs.clustersNative->clustersLinear[clusters[ihit].num]; + const float clamp = cl.getQtot(); + xx *= clamp; + yy *= clamp; + zz *= clamp; + clusterState = clusters[ihit].state; + count = clamp; + } else { + xx = yy = zz = count = 0.f; + clusterState = 0; + } + do { + ihit += wayDirection; + const ClusterNative& GPUrestrict() cl = merger.GetConstantMem()->ioPtrs.clustersNative->clustersLinear[clusters[ihit].num]; + const float clamp = cl.getQtot(); + float clx, cly, clz; + merger.GetConstantMem()->calibObjects.fastTransform->Transform(clusters[ihit].sector, clusters[ihit].row, cl.getPad(), cl.getTime(), clx, cly, clz, mTOffset); + if (chkFunction(ihit, cly, clz)) { + xx += clx * clamp; yy += cly * clamp; zz += clz * clamp; clusterState |= clusters[ihit].state; count += clamp; } - if (!(ihit + wayDirection >= 0 && ihit + wayDirection < maxN && clusters[ihit].row == clusters[ihit + wayDirection].row && clusters[ihit].sector == clusters[ihit + wayDirection].sector)) { - break; - } - ihit += wayDirection; - } + } while (ihit + wayDirection >= 0 && ihit + wayDirection < maxN && clusters[ihit].row == clusters[ihit + wayDirection].row && clusters[ihit].sector == clusters[ihit + wayDirection].sector); if (count < 0.1f) { CADEBUG(printf("\t\tNo matching cluster in double-row, skipping\n")); + xx = tmpX; return -1; } xx /= count; yy /= count; zz /= count; } + if (merger.Param().rec.tpc.rejectIFCLowRadiusCluster) { + const float r2 = xx * xx + yy * yy; + const float rmax2 = CAMath::Square(83.5f + merger.Param().rec.tpc.sysClusErrorMinDist); + if (r2 < rmax2) { + if (markReject) { + MarkClusters(clusters, ihitFirst, ihit, wayDirection, GPUTPCGMMergedTrackHit::flagRejectErr); + } + return -1; + } + } return 0; } -GPUd() float GPUTPCGMTrackParam::AttachClusters(const GPUTPCGMMerger* GPUrestrict() Merger, int32_t sector, int32_t iRow, int32_t iTrack, bool goodLeg, GPUTPCGMPropagator& prop) +GPUd() float GPUTPCGMTrackParam::AttachClusters(const GPUTPCGMMerger& GPUrestrict() merger, int32_t sector, int32_t iRow, int32_t iTrack, bool goodLeg, GPUTPCGMPropagator& prop) { float Y, Z; float X = 0; - Merger->GetConstantMem()->calibObjects.fastTransform->InverseTransformYZtoX(sector, iRow, mP[0], mP[1], X); + merger.GetConstantMem()->calibObjects.fastTransform->InverseTransformYZtoX(sector, iRow, mP[0], mP[1], X); if (prop.GetPropagatedYZ(X, Y, Z)) { Y = mP[0]; Z = mP[1]; } - return AttachClusters(Merger, sector, iRow, iTrack, goodLeg, Y, Z); + return AttachClusters(merger, sector, iRow, iTrack, goodLeg, Y, Z); } -GPUd() float GPUTPCGMTrackParam::AttachClusters(const GPUTPCGMMerger* GPUrestrict() Merger, int32_t sector, int32_t iRow, int32_t iTrack, bool goodLeg, float Y, float Z) +GPUd() float GPUTPCGMTrackParam::AttachClusters(const GPUTPCGMMerger& GPUrestrict() merger, int32_t sector, int32_t iRow, int32_t iTrack, bool goodLeg, float Y, float Z) { - const auto& param = Merger->Param(); - if (param.rec.tpc.disableRefitAttachment & 1) { + const GPUParam& GPUrestrict() param = merger.Param(); + if (param.rec.tpc.disableMarkAdjacent & 1) { return -1e6f; } - const GPUTPCTracker& GPUrestrict() tracker = *(Merger->GetConstantMem()->tpcTrackers + sector); + const GPUTPCTracker& GPUrestrict() tracker = *(merger.GetConstantMem()->tpcTrackers + sector); const GPUTPCRow& GPUrestrict() row = tracker.Row(iRow); GPUglobalref() const cahit2* hits = tracker.HitData(row); GPUglobalref() const calink* firsthit = tracker.FirstHitInBin(row); @@ -501,7 +1018,7 @@ GPUd() float GPUTPCGMTrackParam::AttachClusters(const GPUTPCGMMerger* GPUrestric return -1e6f; } - const float zOffset = param.par.continuousTracking ? Merger->GetConstantMem()->calibObjects.fastTransform->convVertexTimeToZOffset(sector, mTOffset, param.continuousMaxTimeBin) : 0; // TODO: do some validatiomns for the transform conv functions... + const float zOffset = param.par.continuousTracking ? merger.GetConstantMem()->calibObjects.fastTransform->convVertexTimeToZOffset(sector, mTOffset, param.continuousMaxTimeBin) : 0; // TODO: do some validatiomns for the transform conv functions... const float y0 = row.Grid().YMin(); const float stepY = row.HstepY(); const float z0 = row.Grid().ZMin() - zOffset; // We can use our own ZOffset, since this is only used temporarily anyway @@ -509,7 +1026,7 @@ GPUd() float GPUTPCGMTrackParam::AttachClusters(const GPUTPCGMMerger* GPUrestric int32_t bin, ny, nz; float uncorrectedY, uncorrectedZ; - Merger->GetConstantMem()->calibObjects.fastTransform->InverseTransformYZtoNominalYZ(sector, iRow, Y, Z, uncorrectedY, uncorrectedZ); + merger.GetConstantMem()->calibObjects.fastTransform->InverseTransformYZtoNominalYZ(sector, iRow, Y, Z, uncorrectedY, uncorrectedZ); if (CAMath::Abs(uncorrectedY) > row.getTPCMaxY()) { return uncorrectedY; } @@ -521,7 +1038,7 @@ GPUd() float GPUTPCGMTrackParam::AttachClusters(const GPUTPCGMMerger* GPUrestric const float tubeMinSize2 = protect ? param.rec.tpc.tubeProtectMinSize2 : 0.f; float tubeSigma2 = protect ? param.rec.tpc.tubeProtectSigma2 : param.rec.tpc.tubeRemoveSigma2; uint32_t pad = CAMath::Float2UIntRn(GPUTPCGeometry::LinearY2Pad(sector, iRow, uncorrectedY)); - float time = Merger->GetConstantMem()->calibObjects.fastTransform->InverseTransformInTimeFrame(sector, uncorrectedZ + (param.par.continuousTracking ? Merger->GetConstantMem()->calibObjects.fastTransform->convVertexTimeToZOffset(sector, mTOffset, param.continuousMaxTimeBin) : 0), param.continuousMaxTimeBin); // TODO: Simplify this call in TPCFastTransform + float time = merger.GetConstantMem()->calibObjects.fastTransform->InverseTransformInTimeFrame(sector, uncorrectedZ + (param.par.continuousTracking ? merger.GetConstantMem()->calibObjects.fastTransform->convVertexTimeToZOffset(sector, mTOffset, param.continuousMaxTimeBin) : 0), param.continuousMaxTimeBin); // TODO: Simplify this call in TPCFastTransform if (iRow < param.rec.tpc.tubeExtraProtectMinRow || pad < param.rec.tpc.tubeExtraProtectEdgePads || pad >= (uint32_t)(GPUTPCGeometry::NPads(iRow) - param.rec.tpc.tubeExtraProtectEdgePads) || param.GetUnscaledMult(time) / GPUTPCGeometry::Row2X(iRow) > param.rec.tpc.tubeExtraProtectMinOccupancy) { @@ -538,8 +1055,8 @@ GPUd() float GPUTPCGMTrackParam::AttachClusters(const GPUTPCGMMerger* GPUrestric const int32_t nBinsY = row.Grid().Ny(); const int32_t idOffset = tracker.Data().ClusterIdOffset(); const int32_t* ids = &(tracker.Data().ClusterDataIndex()[row.HitNumberOffset()]); - uint32_t myWeight = Merger->TrackOrderAttach()[iTrack] | gputpcgmmergertypes::attachAttached | gputpcgmmergertypes::attachTube; - GPUAtomic(uint32_t)* const weights = Merger->ClusterAttachment(); + uint32_t myWeight = merger.TrackOrderAttach()[iTrack] | gputpcgmmergertypes::attachAttached | gputpcgmmergertypes::attachTube; + GPUAtomic(uint32_t)* const weights = merger.ClusterAttachment(); if (goodLeg) { myWeight |= gputpcgmmergertypes::attachGoodLeg; } @@ -573,42 +1090,6 @@ GPUd() float GPUTPCGMTrackParam::AttachClusters(const GPUTPCGMMerger* GPUrestric return uncorrectedY; } -GPUd() bool GPUTPCGMTrackParam::AttachClustersPropagate(const GPUTPCGMMerger* GPUrestrict() Merger, int32_t sector, int32_t lastRow, int32_t toRow, int32_t iTrack, bool goodLeg, GPUTPCGMPropagator& GPUrestrict() prop, bool inFlyDirection, float maxSinPhi, bool dodEdx) -{ - constexpr float kSectAngle = 2 * M_PI / 18.f; - if (Merger->Param().rec.tpc.disableRefitAttachment & 2) { - return dodEdx; - } - if (CAMath::Abs(lastRow - toRow) < 2) { - return dodEdx; - } - int32_t step = toRow > lastRow ? 1 : -1; - float xx = mX - GPUTPCGeometry::Row2X(lastRow); - for (int32_t iRow = lastRow + step; iRow != toRow; iRow += step) { - if (CAMath::Abs(mP[2]) > maxSinPhi) { - return dodEdx; - } - if (CAMath::Abs(mP[0]) > CAMath::Abs(mX) * CAMath::Tan(kSectAngle / 2.f)) { - return dodEdx; - } - int32_t err = prop.PropagateToXAlpha(xx + GPUTPCGeometry::Row2X(iRow), prop.GetAlpha(), inFlyDirection); - if (err) { - return dodEdx; - } - if (dodEdx && iRow + step == toRow) { - float yUncorrected, zUncorrected; - Merger->GetConstantMem()->calibObjects.fastTransform->InverseTransformYZtoNominalYZ(sector, iRow, mP[0], mP[1], yUncorrected, zUncorrected); - uint32_t pad = CAMath::Float2UIntRn(GPUTPCGeometry::LinearY2Pad(sector, iRow, yUncorrected)); - if (pad >= GPUTPCGeometry::NPads(iRow) || (Merger->GetConstantMem()->calibObjects.dEdxCalibContainer && Merger->GetConstantMem()->calibObjects.dEdxCalibContainer->isDead(sector, iRow, pad))) { - dodEdx = false; - } - } - CADEBUG(printf("Attaching in row %d\n", iRow)); - AttachClusters(Merger, sector, iRow, iTrack, goodLeg, prop); - } - return dodEdx; -} - GPUdii() void GPUTPCGMTrackParam::StoreOuter(gputpcgmmergertypes::GPUTPCOuterParam* outerParam, float alpha) { CADEBUG(printf("\t%21sStorO Alpha %8.3f , X %8.3f - Y %8.3f, Z %8.3f - QPt %7.2f, SP %5.2f --- Cov sY %8.3f sZ %8.3f sSP %8.3f sPt %8.3f\n", "", alpha, mX, mP[0], mP[1], mP[4], mP[2], sqrtf(mC[0]), sqrtf(mC[2]), sqrtf(mC[5]), sqrtf(mC[14]))); @@ -622,7 +1103,7 @@ GPUdii() void GPUTPCGMTrackParam::StoreOuter(gputpcgmmergertypes::GPUTPCOuterPar outerParam->alpha = alpha; } -GPUdic(0, 1) void GPUTPCGMTrackParam::StoreLoopPropagation(const GPUTPCGMMerger* GPUrestrict() Merger, int32_t sector, int32_t iRow, int32_t iTrack, bool outwards, float alpha) +GPUdic(0, 1) void GPUTPCGMTrackParam::StoreLoopPropagation(const GPUTPCGMMerger& GPUrestrict() merger, int32_t sector, int32_t iRow, int32_t iTrack, bool outwards, float alpha) { if (iRow == 0 || iRow == GPUTPCGeometry::NROWS - 1) { return; @@ -637,10 +1118,10 @@ GPUdic(0, 1) void GPUTPCGMTrackParam::StoreLoopPropagation(const GPUTPCGMMerger* return; } - uint32_t nLoopData = CAMath::AtomicAdd(&Merger->Memory()->nLoopData, 1u); - if (nLoopData >= Merger->NMaxTracks()) { - Merger->raiseError(GPUErrors::ERROR_MERGER_LOOPER_OVERFLOW, nLoopData, Merger->NMaxTracks()); - CAMath::AtomicExch(&Merger->Memory()->nLoopData, Merger->NMaxTracks()); + uint32_t nLoopData = CAMath::AtomicAdd(&merger.Memory()->nLoopData, 1u); + if (nLoopData >= merger.NMaxTracks()) { + merger.raiseError(GPUErrors::ERROR_MERGER_LOOPER_OVERFLOW, nLoopData, merger.NMaxTracks()); + CAMath::AtomicExch(&merger.Memory()->nLoopData, merger.NMaxTracks()); return; } GPUTPCGMLoopData data; @@ -650,36 +1131,38 @@ GPUdic(0, 1) void GPUTPCGMTrackParam::StoreLoopPropagation(const GPUTPCGMMerger* data.sector = sector; data.row = iRow; data.outwards = outwards; - Merger->LoopData()[nLoopData] = data; + merger.LoopData()[nLoopData] = data; } -GPUdii() void GPUTPCGMTrackParam::PropagateLooper(const GPUTPCGMMerger* GPUrestrict() Merger, int32_t loopIdx) +GPUdii() void GPUTPCGMTrackParam::PropagateLooper(const GPUTPCGMMerger& GPUrestrict() merger, int32_t loopIdx) { + GPUTPCGMLoopData& data = merger.LoopData()[loopIdx]; + if (!merger.MergedTracks()[data.track].OK()) { + return; + } GPUTPCGMPropagator prop; prop.SetMaterialTPC(); - prop.SetPolynomialField(&Merger->Param().polynomialField); + prop.SetPolynomialField(&merger.Param().polynomialField); prop.SetMaxSinPhi(constants::MAX_SIN_PHI); - prop.SetMatLUT(Merger->Param().rec.useMatLUT ? Merger->GetConstantMem()->calibObjects.matLUT : nullptr); + prop.SetMatLUT(merger.Param().rec.useMatLUT ? merger.GetConstantMem()->calibObjects.matLUT : nullptr); prop.SetSeedingErrors(false); prop.SetFitInProjections(true); prop.SetPropagateBzOnly(false); - GPUTPCGMLoopData& data = Merger->LoopData()[loopIdx]; prop.SetTrack(&data.param, data.alpha); - if (false) { - data.param.AttachClustersLooper(Merger, data.sector, data.row, data.track, data.outwards, prop); + if (merger.Param().rec.tpc.looperFollowMode == 1) { + data.param.AttachClustersLooperFollow(merger, prop, data.sector, data.track, data.outwards); } else { - data.param.AttachClustersLooperFollow(Merger, prop, data.sector, data.track, data.outwards); + data.param.AttachClustersLooper(merger, data.sector, data.row, data.track, data.outwards, prop); } } -GPUdi() void GPUTPCGMTrackParam::AttachClustersLooperFollow(const GPUTPCGMMerger* GPUrestrict() Merger, GPUTPCGMPropagator& GPUrestrict() prop, int32_t sector, int32_t iTrack, bool up) +GPUdi() void GPUTPCGMTrackParam::AttachClustersLooperFollow(const GPUTPCGMMerger& GPUrestrict() merger, GPUTPCGMPropagator& GPUrestrict() prop, int32_t sector, const int32_t iTrack, const bool up) { float toX = mX; - bool inFlyDirection = (Merger->MergedTracks()[iTrack].Leg() & 1) ^ up; + bool inFlyDirection = (merger.MergedTracks()[iTrack].Leg() & 1) ^ up; - constexpr float kSectAngle = 2 * M_PI / 18.f; - const GPUParam& GPUrestrict() param = Merger->Param(); + const GPUParam& GPUrestrict() param = merger.Param(); bool right = (mP[2] < 0) ^ up; const int32_t sectorSide = sector >= (int32_t)(GPUTPCGeometry::NSECTORS / 2) ? (GPUTPCGeometry::NSECTORS / 2) : 0; float lrFactor = right ^ !up ? 1.f : -1.f; @@ -693,7 +1176,7 @@ GPUdi() void GPUTPCGMTrackParam::AttachClustersLooperFollow(const GPUTPCGMMerger CADEBUG(printf("\tRotated: X %f Y %f Z %f SinPhi %f (Alpha %f / %f)\n", mP[0], mX, mP[1], mP[2], prop.GetAlpha(), prop.GetAlpha() + CAMath::Pi() / 2.f)); uint32_t maxTries = 100; while (true) { - while (CAMath::Abs(mX) <= CAMath::Abs(mP[0]) * CAMath::Tan(kSectAngle / 2.f) + 0.1f) { + while (CAMath::Abs(mX) <= CAMath::Abs(mP[0]) * CAMath::Tan(GPUTPCGeometry::kSectAngle() / 2.f) + 0.1f) { if (maxTries-- == 0) { return; } @@ -716,7 +1199,7 @@ GPUdi() void GPUTPCGMTrackParam::AttachClustersLooperFollow(const GPUTPCGMMerger float rowX = GPUTPCGeometry::Row2X(j); if (CAMath::Abs(rowX - (-mP[0] * lrFactor)) < 1.5f) { CADEBUG(printf("\t\tAttempt row %d (X %f Y %f Z %f)\n", j, rowX, mX * lrFactor, mP[1])); - AttachClusters(Merger, sector, j, iTrack, false, mX * lrFactor, mP[1]); + AttachClusters(merger, sector, j, iTrack, false, mX * lrFactor, mP[1]); } } } @@ -742,9 +1225,8 @@ GPUdi() void GPUTPCGMTrackParam::AttachClustersLooperFollow(const GPUTPCGMMerger } } -GPUdi() void GPUTPCGMTrackParam::AttachClustersLooper(const GPUTPCGMMerger* GPUrestrict() Merger, int32_t sector, int32_t iRow, int32_t iTrack, bool outwards, GPUTPCGMPropagator& GPUrestrict() prop) +GPUdi() void GPUTPCGMTrackParam::AttachClustersLooper(const GPUTPCGMMerger& GPUrestrict() merger, int32_t sector, int32_t iRow, const int32_t iTrack, const bool up, const GPUTPCGMPropagator& GPUrestrict() prop) { - constexpr float kSectAngle = 2 * M_PI / 18.f; // Note that the coordinate system is rotated by 90 degree swapping X and Y! float X = mP[2] > 0 ? mP[0] : -mP[0]; float Y = mP[2] > 0 ? -mX : mX; @@ -752,7 +1234,7 @@ GPUdi() void GPUTPCGMTrackParam::AttachClustersLooper(const GPUTPCGMMerger* GPUr float SinPhi = CAMath::Sqrt(1 - mP[2] * mP[2]) * (mP[2] > 0 ? -1 : 1); float b = prop.GetBz(prop.GetAlpha(), mX, mP[0], mP[1]); - float dx = outwards ? 1.f : -1.f; + float dx = up ? 1.f : -1.f; const float myRowX = GPUTPCGeometry::Row2X(iRow); // printf("\nAttachMirror sector %d row %d outwards %d\n", (int)sector, (int)iRow, (int)outwards); // printf("X %f Y %f Z %f SinPhi %f -->\n", mX, mP[0], mP[1], mP[2]); @@ -779,44 +1261,45 @@ GPUdi() void GPUTPCGMTrackParam::AttachClustersLooper(const GPUTPCGMMerger* GPUr Y += dS * SinPhi + h4; Z += dS * mP[3]; SinPhi = newSinPhi; - if (CAMath::Abs(X) > CAMath::Abs(Y) * CAMath::Tan(kSectAngle / 2.f)) { + if (CAMath::Abs(X) > CAMath::Abs(Y) * CAMath::Tan(GPUTPCGeometry::kSectAngle() / 2.f)) { // printf("Abort, sector edge\n"); return; } // printf("count %d: At X %f Y %f Z %f SinPhi %f\n", maxTries, mP[2] > 0 ? -Y : Y, mP[2] > 0 ? X : -X, Z, SinPhi); float paramX = mP[2] > 0 ? -Y : Y; - int32_t step = outwards ? 1 : -1; + int32_t step = up ? 1 : -1; int32_t found = 0; for (int32_t j = iRow; j >= 0 && j < (int32_t)GPUTPCGeometry::NROWS && found < 3; j += step) { float rowX = mX + GPUTPCGeometry::Row2X(j) - myRowX; if (CAMath::Abs(rowX - paramX) < 1.5f) { // printf("Attempt row %d at y %f\n", j, X); - AttachClusters(Merger, sector, j, iTrack, false, mP[2] > 0 ? X : -X, Z); + AttachClusters(merger, sector, j, iTrack, false, mP[2] > 0 ? X : -X, Z); } } } } -GPUd() float GPUTPCGMTrackParam::ShiftZ(const GPUTPCGMMergedTrackHit* clusters, const GPUTPCGMMerger* merger, int32_t N) +GPUd() float GPUTPCGMTrackParam::ShiftZ(const GPUTPCGMMergedTrackHit* GPUrestrict() clusters, const GPUTPCGMMerger& GPUrestrict() merger, int32_t N) { if (N == 0) { N = 1; } - const auto& GPUrestrict() cls = merger->GetConstantMem()->ioPtrs.clustersNative->clustersLinear; + const auto& GPUrestrict() cls = merger.GetConstantMem()->ioPtrs.clustersNative->clustersLinear; float z0 = cls[clusters[0].num].getTime(), zn = cls[clusters[N - 1].num].getTime(); const auto tmp = zn > z0 ? std::array{zn, z0, GPUTPCGeometry::Row2X(clusters[N - 1].row)} : std::array{z0, zn, GPUTPCGeometry::Row2X(clusters[0].row)}; return ShiftZ(merger, clusters[0].sector, tmp[0], tmp[1], tmp[2]); } -GPUd() float GPUTPCGMTrackParam::ShiftZ(const GPUTPCGMMerger* GPUrestrict() merger, uint32_t sector, float cltmax, float cltmin, float clx) +GPUd() float GPUTPCGMTrackParam::ShiftZ(const GPUTPCGMMerger& GPUrestrict() merger, uint32_t sector, float cltmax, float cltmin, float clx) { - if (!merger->Param().par.continuousTracking) { + const GPUParam& GPUrestrict() param = merger.Param(); + if (!param.par.continuousTracking) { return 0.f; } float deltaZ = 0.f; bool beamlineReached = false; - const float r1 = CAMath::Max(0.0001f, CAMath::Abs(mP[4] * merger->Param().polynomialField.GetNominalBz())); + const float r1 = CAMath::Max(0.0001f, CAMath::Abs(mP[4] * param.polynomialField.GetNominalBz())); if (r1 < 0.01501) { // 100 MeV @ 0.5T ~ 0.66m cutof const float dist2 = mX * mX + mP[0] * mP[0]; const float dist1r2 = dist2 * r1 * r1; @@ -843,33 +1326,34 @@ GPUd() float GPUTPCGMTrackParam::ShiftZ(const GPUTPCGMMerger* GPUrestrict() merg } if (!beamlineReached) { - float refZ = ((sector < GPUTPCGeometry::NSECTORS / 2) ? merger->Param().rec.tpc.defaultZOffsetOverR : -merger->Param().rec.tpc.defaultZOffsetOverR) * clx; + float refZ = ((sector < GPUTPCGeometry::NSECTORS / 2) ? param.rec.tpc.defaultZOffsetOverR : -param.rec.tpc.defaultZOffsetOverR) * clx; float basez; - merger->GetConstantMem()->calibObjects.fastTransform->TransformIdealZ(sector, cltmax, basez, mTOffset); + merger.GetConstantMem()->calibObjects.fastTransform->TransformIdealZ(sector, cltmax, basez, mTOffset); deltaZ = basez - refZ; } - { - float deltaT = merger->GetConstantMem()->calibObjects.fastTransform->convDeltaZtoDeltaTimeInTimeFrame(sector, deltaZ); + mTOffset += merger.GetConstantMem()->calibObjects.fastTransform->convDeltaZtoDeltaTimeInTimeFrame(sector, deltaZ); + return -ConstrainZ(merger, sector, cltmax, cltmin, deltaZ); +} + +GPUd() float GPUTPCGMTrackParam::ConstrainZ(const GPUTPCGMMerger& GPUrestrict() merger, uint32_t sector, float cltmax, float cltmin, float deltaZ) +{ + const float maxT = cltmin - merger.GetConstantMem()->calibObjects.fastTransform->getT0(); + const float minT = cltmax - merger.GetConstantMem()->calibObjects.fastTransform->getMaxDriftTime(sector); + // printf("T Check: Clusters %f %f, min %f max %f vtx %f\n", tz1, tz2, minT, maxT, mTOffset); + float deltaT = 0.f; + if (mTOffset < minT) { + deltaT = minT - mTOffset; + } + if (mTOffset + deltaT > maxT) { + deltaT = maxT - mTOffset; + } + if (deltaT != 0.f) { + deltaZ += merger.GetConstantMem()->calibObjects.fastTransform->convDeltaTimeToDeltaZinTimeFrame(sector, deltaT); + // printf("Moving clusters to TPC Range: QPt %f, New mTOffset %f, t1 %f, t2 %f, Shift %f in Z: %f to %f --> %f to %f in T\n", mP[4], mTOffset + deltaT, tz1, tz2, deltaZ, tz2 - mTOffset, tz1 - mTOffset, tz2 - mTOffset - deltaT, tz1 - mTOffset - deltaT); mTOffset += deltaT; - const float maxT = cltmin - merger->GetConstantMem()->calibObjects.fastTransform->getT0(); - const float minT = cltmax - merger->GetConstantMem()->calibObjects.fastTransform->getMaxDriftTime(sector); - // printf("T Check: Clusters %f %f, min %f max %f vtx %f\n", tz1, tz2, minT, maxT, mTOffset); - deltaT = 0.f; - if (mTOffset < minT) { - deltaT = minT - mTOffset; - } - if (mTOffset + deltaT > maxT) { - deltaT = maxT - mTOffset; - } - if (deltaT != 0.f) { - deltaZ += merger->GetConstantMem()->calibObjects.fastTransform->convDeltaTimeToDeltaZinTimeFrame(sector, deltaT); - // printf("Moving clusters to TPC Range: QPt %f, New mTOffset %f, t1 %f, t2 %f, Shift %f in Z: %f to %f --> %f to %f in T\n", mP[4], mTOffset + deltaT, tz1, tz2, deltaZ, tz2 - mTOffset, tz1 - mTOffset, tz2 - mTOffset - deltaT, tz1 - mTOffset - deltaT); - mTOffset += deltaT; - } - mP[1] -= deltaZ; } - // printf("\n"); - return -deltaZ; + mP[1] -= deltaZ; + return deltaZ; } GPUd() bool GPUTPCGMTrackParam::CheckCov() const @@ -904,54 +1388,52 @@ GPUd() bool GPUTPCGMTrackParam::CheckNumericalQuality(float overrideCovYY) const return ok; } -GPUdii() void GPUTPCGMTrackParam::RefitTrack(GPUTPCGMMergedTrack& GPUrestrict() track, int32_t iTrk, GPUTPCGMMerger* GPUrestrict() merger, int32_t attempt) // VS: GPUd changed to GPUdii. No change in output and no performance penalty. +GPUdii() void GPUTPCGMTrackParam::RefitTrack(GPUTPCGMMergedTrack& GPUrestrict() track, int32_t iTrk, GPUTPCGMMerger& GPUrestrict() merger, bool rebuilt, bool retryAttempt) { if (!track.OK()) { return; } - // clang-format off - CADEBUG(if (DEBUG_SINGLE_TRACK != -1 && iTrk != ((DEBUG_SINGLE_TRACK == -2 && getenv("DEBUG_TRACK")) ? atoi(getenv("DEBUG_TRACK")) : DEBUG_SINGLE_TRACK)) { track.SetNClusters(0); track.SetOK(0); return; } ); - // clang-format on + CADEBUG3(if (GPUCA_CADEBUG_ENABLED >= 2) { static const std::set dbgTracks = []() { std::set m; if (const char* e = getenv("DEBUG_TRACK")) { std::string s(e); for (char* t = strtok(s.data(), ";, \n\t"); t; t = strtok(nullptr, ";, \n\t")) { m.insert(atoi(t)); } } return m; }(); CADEBUG_CONDITION = dbgTracks.count(iTrk); }); int32_t nTrackHits = track.NClusters(); int32_t NTolerated = 0; // Clusters not fit but tollerated for track length cut GPUTPCGMTrackParam t = track.Param(); float Alpha = track.Alpha(); - CADEBUG(int32_t nTrackHitsOld = nTrackHits; float ptOld = t.QPt()); - bool ok = t.Fit(merger, iTrk, merger->Clusters() + track.FirstClusterRef(), nTrackHits, NTolerated, Alpha, attempt, constants::MAX_SIN_PHI, track); - 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(), (int32_t)ok, t.Chi2(), t.Chi2() / CAMath::Max(1, nTrackHits))); + bool ok = t.Fit(merger, iTrk, nTrackHits, NTolerated, Alpha, track, rebuilt, retryAttempt); + CADEBUG(if (!merger.Param().rec.tpc.rebuildTrackInFit || rebuilt) printf("Finished Fit Track %7d --- OUTPUT hits %d -> %d+%d = %d, QPt %f -> %f, SP %f, OK %d chi2 %f chi2ndf %f\n", iTrk, track.NClusters(), nTrackHits, NTolerated, nTrackHits + NTolerated, track.GetParam().GetQPt(), t.QPt(), t.SinPhi(), (int32_t)ok, t.Chi2(), t.Chi2() / CAMath::Max(1, nTrackHits))); - if (!ok && attempt == 0 && merger->Param().rec.tpc.retryRefit) { + if (!ok && (!merger.Param().rec.tpc.rebuildTrackInFit || rebuilt) && !retryAttempt && merger.Param().rec.tpc.retryRefit) { for (uint32_t i = 0; i < track.NClusters(); i++) { - merger->Clusters()[track.FirstClusterRef() + i].state &= GPUTPCGMMergedTrackHit::clustererAndSharedFlags; + merger.Clusters()[track.FirstClusterRef() + i].state &= GPUTPCGMMergedTrackHit::clustererAndSharedFlags; } CADEBUG(printf("Track rejected, marking for retry\n")); - if (merger->Param().rec.tpc.retryRefit == 2) { - nTrackHits = track.NClusters(); - 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(), nTrackHits, NTolerated, Alpha, 1, constants::MAX_SIN_PHI, track); - } else { - uint32_t nRefit = CAMath::AtomicAdd(&merger->Memory()->nRetryRefit, 1u); - merger->RetryRefitIds()[nRefit] = iTrk; - return; - } + uint32_t nRefit = CAMath::AtomicAdd(&merger.Memory()->nRetryRefit, 1u); + merger.RetryRefitIds()[nRefit] = iTrk; // TODO: With rebuildTrackInFit the retry runs on the rebuilt clusters, so it cannot recover tracks that lost their clusters in the rebuild + return; + } + if (retryAttempt && (t.mNDF < 10 || t.mChi2 / t.mNDF >= 6.f)) { + ok = false; } + if (CAMath::Abs(t.QPt()) < 1.e-4f) { - t.QPt() = 1.e-4f; + t.QPt() = CAMath::Copysign(1.e-4f, t.QPt()); } - CADEBUG(if (t.GetX() > 250) { printf("ERROR, Track %d at impossible X %f, Pt %f, Looper %d\n", iTrk, t.GetX(), CAMath::Abs(1.f / t.QPt()), (int32_t)merger->MergedTracks()[iTrk].Looper()); }); + CADEBUG(if (t.GetX() > 250) { printf("ERROR, Track %d at impossible X %f, Pt %f, Looper %d\n", iTrk, t.GetX(), CAMath::Abs(1.f / t.QPt()), (int32_t)merger.MergedTracks()[iTrk].Looper()); }); track.SetOK(ok); - track.SetNClustersFitted(nTrackHits); - track.Param() = t; - track.Alpha() = Alpha; + if (t.GetNDF() <= 0 && !rebuilt && merger.Param().rec.tpc.rebuildTrackInFit) { // TODO: Better handling of NDF<0 tracks, how do we want to do cluster rejection? + track.Param().NDF() = 0; + } else { + track.Param() = t; + track.Alpha() = Alpha; + } + if (!merger.Param().rec.tpc.rebuildTrackInFit || rebuilt) { + track.SetNClustersFitted(nTrackHits); + } - // if (track.OK()) merger->DebugRefitMergedTrack(track); + // if (track.OK()) merger.DebugRefitMergedTrack(track); } GPUd() void GPUTPCGMTrackParam::Rotate(float alpha) diff --git a/GPU/GPUTracking/Merger/GPUTPCGMTrackParam.h b/GPU/GPUTracking/Merger/GPUTPCGMTrackParam.h index c9d8f2017af52..1d98920cfb094 100644 --- a/GPU/GPUTracking/Merger/GPUTPCGMTrackParam.h +++ b/GPU/GPUTracking/Merger/GPUTPCGMTrackParam.h @@ -27,8 +27,14 @@ class AliExternalTrackParam; +namespace o2::tpc +{ +struct ClusterNative; +}; + namespace o2::gpu { +class GPUTPCTracker; class GPUTPCGMMerger; class GPUTPCGMBorderTrack; struct GPUParam; @@ -36,6 +42,13 @@ class GPUTPCGMPhysicalTrackModel; class GPUTPCGMPolynomialField; class GPUTPCGMMergedTrack; class GPUTPCGMPropagator; +class GPUdEdx; +struct GPUTPCGMMergedTrackHit; + +namespace gputpcgmmergertypes +{ +struct InterpolationErrorHit; +} // namespace gputpcgmmergertypes /** * @class GPUTPCGMTrackParam @@ -141,22 +154,28 @@ class GPUTPCGMTrackParam GPUd() bool CheckNumericalQuality(float overrideCovYY = -1.f) const; GPUd() bool CheckCov() const; - GPUd() bool Fit(GPUTPCGMMerger* merger, int32_t iTrk, GPUTPCGMMergedTrackHit* clusters, int32_t& N, int32_t& NTolerated, float& Alpha, int32_t attempt, float maxSinPhi, GPUTPCGMMergedTrack& track); + GPUd() bool Fit(GPUTPCGMMerger& merger, int32_t iTrk, int32_t& N, int32_t& NTolerated, float& Alpha, GPUTPCGMMergedTrack& track, bool rebuilt, bool retryAttempt); + GPUd() void DodEdx(GPUdEdx& dEdx, GPUdEdx& dEdxAlt, GPUTPCGMMerger& merger, bool finalFit, int ihit, int ihitMergeFirst, int wayDirection, const GPUTPCGMMergedTrackHit* clusters, uint8_t clusterState, float zz, uint8_t dEdxSubThresholdRow GPUCA_DEBUG_STREAMER_CHECK(, int32_t iTrk)); + GPUd() int32_t FitHit(GPUTPCGMMerger& merger, const int32_t iTrk, const GPUTPCGMMergedTrack& track, const float xx, const float yy, const float zz, const uint8_t clusterState, const float clAlpha, const int32_t iWay, const bool inFlyDirection, float& deltaZ, float& lastUpdateX, GPUTPCGMMergedTrackHit* clusters, GPUTPCGMPropagator& prop, gputpcgmmergertypes::InterpolationErrorHit& inter, GPUdEdx& dEdx, GPUdEdx& dEdxAlt, float& sumInvSqrtCharge, int32_t& nAvgCharge, const int32_t ihit, const int32_t ihitMergeFirst, const bool allowChangeClusters, const bool refit, const bool finalFit, int32_t& nMissed, int32_t& nMissed2, float uncorrectedY, bool retryAttempt); + GPUd() void FitAddRow(const int32_t iRow, const uint8_t sector, const int32_t iTrk, const GPUTPCGMMergedTrack& track, GPUTPCGMPropagator& prop, const bool inFlyDirection, GPUTPCGMMerger& merger, uint8_t* dEdxSubThresholdRow, const bool dodEdx, const bool doAttach, const uint8_t doInterpolate, gputpcgmmergertypes::InterpolationErrorHit& inter, const float deltaZ, const float sumInvSqrtCharge, const int32_t nAvgCharge); + GPUd() void HandleCrossCE(const GPUParam& param, const uint8_t sector, const uint8_t& lastSector); + GPUd() static void RefitTrack(GPUTPCGMMergedTrack& track, int32_t iTrk, GPUTPCGMMerger& merger, bool rebuilt, bool retryAttempt); GPUd() void MoveToReference(GPUTPCGMPropagator& prop, const GPUParam& param, float& alpha); GPUd() void MirrorTo(GPUTPCGMPropagator& prop, float toY, float toZ, bool inFlyDirection, const GPUParam& param, uint8_t row, uint8_t clusterState, bool mirrorParameters, int8_t sector); - GPUd() int32_t MergeDoubleRowClusters(int32_t& ihit, int32_t wayDirection, GPUTPCGMMergedTrackHit* clusters, const GPUTPCGMMerger* merger, GPUTPCGMPropagator& prop, float& xx, float& yy, float& zz, int32_t maxN, float clAlpha, uint8_t& clusterState, bool rejectChi2); - - GPUd() bool AttachClustersPropagate(const GPUTPCGMMerger* GPUrestrict() Merger, int32_t sector, int32_t lastRow, int32_t toRow, int32_t iTrack, bool goodLeg, GPUTPCGMPropagator& prop, bool inFlyDirection, float maxSinPhi = constants::MAX_SIN_PHI, bool checkdEdx = false); - GPUd() float AttachClusters(const GPUTPCGMMerger* GPUrestrict() Merger, int32_t sector, int32_t iRow, int32_t iTrack, bool goodLeg, GPUTPCGMPropagator& prop); // Returns uncorrectedY for later use - GPUd() float AttachClusters(const GPUTPCGMMerger* GPUrestrict() Merger, int32_t sector, int32_t iRow, int32_t iTrack, bool goodLeg, float Y, float Z); - GPUd() void AttachClustersLooper(const GPUTPCGMMerger* GPUrestrict() Merger, int32_t sector, int32_t iRow, int32_t iTrack, bool outwards, GPUTPCGMPropagator& prop); - GPUd() void AttachClustersLooperFollow(const GPUTPCGMMerger* GPUrestrict() Merger, GPUTPCGMPropagator& prop, int32_t sector, int32_t iTrack, bool outwards); - GPUd() void StoreLoopPropagation(const GPUTPCGMMerger* GPUrestrict() Merger, int32_t sector, int32_t iRow, int32_t iTrack, bool outwards, float alpha); + GPUd() int32_t MergeDoubleRowClusters(int32_t& ihit, int32_t wayDirection, GPUTPCGMMergedTrackHit* clusters, const GPUTPCGMMerger& merger, GPUTPCGMPropagator& prop, float& xx, float& yy, float& zz, int32_t maxN, float clAlpha, uint8_t& clusterState, const bool markReject); + GPUd() float FindBestInterpolatedHit(GPUTPCGMMerger& merger, gputpcgmmergertypes::InterpolationErrorHit& inter, const uint8_t sector, const uint8_t row, const float deltaZ, const float sumInvSqrtCharge, const int nAvgCharge, const GPUTPCGMPropagator& prop, const int32_t iTrk, bool interOnly, bool extrapolation = false); + GPUd() void InterpolateMissingRows(GPUTPCGMMerger& merger, gputpcgmmergertypes::InterpolationErrors& interpolation, GPUTPCGMMergedTrackHit* clusters, int32_t ihit, int32_t interpolationIndex, int32_t lastRow, const float deltaZ, const float sumInvSqrtCharge, const int32_t nAvgCharge, const GPUTPCGMPropagator& prop, const int32_t iTrk); + + GPUd() float AttachClusters(const GPUTPCGMMerger& GPUrestrict() merger, int32_t sector, int32_t iRow, int32_t iTrack, bool goodLeg, GPUTPCGMPropagator& prop); // Returns uncorrectedY for later use + GPUd() float AttachClusters(const GPUTPCGMMerger& GPUrestrict() merger, int32_t sector, int32_t iRow, int32_t iTrack, bool goodLeg, float Y, float Z); + GPUd() void AttachClustersLooper(const GPUTPCGMMerger& GPUrestrict() merger, int32_t sector, int32_t iRow, const int32_t iTrack, const bool up, const GPUTPCGMPropagator& prop); + GPUd() void AttachClustersLooperFollow(const GPUTPCGMMerger& GPUrestrict() merger, GPUTPCGMPropagator& prop, int32_t sector, const int32_t iTrack, const bool up); + GPUd() void StoreLoopPropagation(const GPUTPCGMMerger& GPUrestrict() merger, int32_t sector, int32_t iRow, int32_t iTrack, bool outwards, float alpha); GPUd() void StoreOuter(gputpcgmmergertypes::GPUTPCOuterParam* outerParam, float alpha); - GPUd() static void PropagateLooper(const GPUTPCGMMerger* GPUrestrict() Merger, int32_t loopIdx); + GPUd() static void PropagateLooper(const GPUTPCGMMerger& merger, int32_t loopIdx); - GPUd() void AddCovDiagErrors(const float* GPUrestrict() errors2); - GPUd() void AddCovDiagErrorsWithCorrelations(const float* GPUrestrict() errors2); + GPUd() void AddCovDiagErrors(const float* errors2); + GPUd() void AddCovDiagErrorsWithCorrelations(const float* errors2); GPUdi() void MarkClusters(GPUTPCGMMergedTrackHit* GPUrestrict() clusters, int32_t ihitFirst, int32_t ihitLast, int32_t wayDirection, uint8_t state) { @@ -182,8 +201,9 @@ class GPUTPCGMTrackParam } GPUd() void Rotate(float alpha); - GPUd() float ShiftZ(const GPUTPCGMMerger* merger, uint32_t sector, float cltmax, float cltmin, float clx); - GPUd() float ShiftZ(const GPUTPCGMMergedTrackHit* clusters, const GPUTPCGMMerger* merger, int32_t N); + GPUd() float ShiftZ(const GPUTPCGMMerger& merger, uint32_t sector, float cltmax, float cltmin, float clx); + GPUd() float ShiftZ(const GPUTPCGMMergedTrackHit* clusters, const GPUTPCGMMerger& merger, int32_t N); + GPUd() float ConstrainZ(const GPUTPCGMMerger& merger, uint32_t sector, float cltmax, float cltmin, float deltaZ = 0.f); // Shift Z / TOffset such that the clusters are within the TPC drift volume, applies in addition the pending shift deltaZ to Z, returns the total shift GPUd() static float Reciprocal(float x) { return 1.f / x; } GPUdi() static void Assign(float& x, bool mask, float v) @@ -200,8 +220,6 @@ class GPUTPCGMTrackParam } } - GPUd() static void RefitTrack(GPUTPCGMMergedTrack& track, int32_t iTrk, GPUTPCGMMerger* merger, int32_t attempt); - GPUdi() void ConstrainSinPhi(float limit = constants::MAX_SIN_PHI) { if (mP[2] > limit) { @@ -212,8 +230,6 @@ class GPUTPCGMTrackParam } private: - GPUd() int32_t initResetT0(); - float mX; // x position float mTOffset; // Z offset with early transform, T offset otherwise float mP[5]; // 'active' track parameters: Y, Z, SinPhi, DzDs, q/Pt @@ -231,15 +247,6 @@ struct GPUTPCGMLoopData { uint8_t outwards; }; -GPUdi() int32_t GPUTPCGMTrackParam::initResetT0() -{ - const float absQPt = CAMath::Abs(mP[4]); - if (absQPt < (150.f / 40.f)) { - return 150.f / 40.f; - } - return CAMath::Max(10.f, 150.f / mP[4]); -} - GPUdi() void GPUTPCGMTrackParam::ResetCovariance() { mC[0] = 100.f; diff --git a/GPU/GPUTracking/Refit/GPUTrackingRefit.cxx b/GPU/GPUTracking/Refit/GPUTrackingRefit.cxx index c8d6b00d916cd..906e943bb1743 100644 --- a/GPU/GPUTracking/Refit/GPUTrackingRefit.cxx +++ b/GPU/GPUTracking/Refit/GPUTrackingRefit.cxx @@ -400,14 +400,11 @@ GPUd() int32_t GPUTrackingRefit::RefitTrack(T& trkX, bool outward, bool resetCov trk.NormalizeAlpha(alpha); prop.SetAlpha(alpha); } else if constexpr (std::is_same_v) { - constexpr float kDeg2Rad = M_PI / 180.f; - constexpr float kSectAngle = 2 * M_PI / 18.f; if (mPparam->rec.tpc.trackReferenceX <= 500) { - // a forced energy-loss sign holds along the refit direction: the way to the reference X can run against it - const int32_t eLossSignRef = (mPparam->rec.tpc.trackReferenceX > trk.getX()) == outward ? eLossSign : -eLossSign; + const int32_t eLossSignRef = (mPparam->rec.tpc.trackReferenceX > trk.getX()) == outward ? eLossSign : -eLossSign; // a forced energy-loss sign holds along the refit direction: the way to the reference X can run against it if (prop->PropagateToXBxByBz(trk, mPparam->rec.tpc.trackReferenceX, Propagator::MAX_SIN_PHI, Propagator::MAX_STEP, Propagator::MatCorrType::USEMatCorrLUT, nullptr, eLossSignRef)) { - if (CAMath::Abs(trk.getY()) > trk.getX() * CAMath::Tan(kSectAngle / 2.f)) { - float newAlpha = trk.getAlpha() + CAMath::Round(CAMath::ATan2(trk.getY(), trk.getX()) / kDeg2Rad / 20.f) * kSectAngle; + if (CAMath::Abs(trk.getY()) > trk.getX() * CAMath::Tan(GPUTPCGeometry::kSectAngle() / 2.f)) { + float newAlpha = trk.getAlpha() + CAMath::Round(CAMath::ATan2(trk.getY(), trk.getX()) / CAMath::Deg2Rad() / 20.f) * GPUTPCGeometry::kSectAngle(); GPUTPCGMTrackParam::NormalizeAlpha(newAlpha); trk.rotate(newAlpha) && prop->PropagateToXBxByBz(trk, mPparam->rec.tpc.trackReferenceX, Propagator::MAX_SIN_PHI, Propagator::MAX_STEP, Propagator::MatCorrType::USEMatCorrLUT, nullptr, eLossSignRef); } diff --git a/GPU/GPUTracking/SectorTracker/GPUTPCTracker.h b/GPU/GPUTracking/SectorTracker/GPUTPCTracker.h index cb0c8d4a76fdb..aaf1bff882522 100644 --- a/GPU/GPUTracking/SectorTracker/GPUTPCTracker.h +++ b/GPU/GPUTracking/SectorTracker/GPUTPCTracker.h @@ -162,11 +162,15 @@ class GPUTPCTracker : public GPUProcessor GPUdi() static int32_t CalculateHitWeight(int32_t NHits, float chi2) { float weight = NHits * (NHits * 2 - 5) * 128 / chi2; // TODO: Add QPt to this formula - if (weight < 0.f || weight > 2e9f) { + if (weight < 0.f) { return 0; } + if (weight > 2e9f) { + return 2e9f; + } return ((int32_t)weight); } + GPUdi() float GetChiSeedFactor() const { return Param().rec.tpc.hitPickUpFactor * Param().rec.tpc.hitPickUpFactor * 3.5f * 3.5f; } GPUd() void MaximizeHitWeight(const GPUTPCRow& row, int32_t hitIndex, int32_t weight) { mData.MaximizeHitWeight(row, hitIndex, weight); } GPUd() void SetHitWeight(const GPUTPCRow& row, int32_t hitIndex, int32_t weight) { mData.SetHitWeight(row, hitIndex, weight); } GPUd() int32_t HitWeight(const GPUTPCRow& row, int32_t hitIndex) const { return mData.HitWeight(row, hitIndex); } diff --git a/GPU/GPUTracking/SectorTracker/GPUTPCTrackletConstructor.cxx b/GPU/GPUTracking/SectorTracker/GPUTPCTrackletConstructor.cxx index 33a3264a87ab3..2f89a998f8e9f 100644 --- a/GPU/GPUTracking/SectorTracker/GPUTPCTrackletConstructor.cxx +++ b/GPU/GPUTracking/SectorTracker/GPUTPCTrackletConstructor.cxx @@ -56,8 +56,8 @@ GPUd() void GPUTPCTrackletConstructor::StoreTracklet(int32_t /*nBlocks*/, int32_ { // reconstruction of tracklets, tracklet store step const uint32_t nHits = r.mLastRow + 1 - r.mFirstRow; - if (nHits == 0 || r.mNHits == 0 || (r.mNHits < tracker.Param().tpcMinHitsB5(tParam.QPt() * tracker.Param().qptB5Scaler) || !CheckCov(tParam) || CAMath::Abs(tParam.GetQPt() * tracker.Param().qptB5Scaler) > tracker.Param().rec.maxTrackQPtB5)) { - CADEBUG(printf(" Rejected: nHits %d QPt %f MinHits %d MaxQPt %f CheckCov %d\n", r.mNHits, tParam.QPt(), tracker.Param().tpcMinHitsB5(tParam.QPt() * tracker.Param().qptB5Scaler), tracker.Param().rec.maxTrackQPtB5, (int32_t)CheckCov(tParam))); + if (nHits == 0 || r.mNHits == 0 || (r.mNHits < tracker.Param().tpcMinHitsB5Abs(CAMath::Abs(tParam.QPt() * tracker.Param().qptB5Scaler)) || !CheckCov(tParam) || CAMath::Abs(tParam.GetQPt() * tracker.Param().qptB5Scaler) > tracker.Param().rec.maxTrackQPtB5)) { + CADEBUG(printf(" Rejected: nHits %d QPt %f MinHits %d MaxQPt %f CheckCov %d\n", r.mNHits, tParam.QPt(), tracker.Param().tpcMinHitsB5Abs(CAMath::Abs(tParam.QPt() * tracker.Param().qptB5Scaler)), tracker.Param().rec.maxTrackQPtB5, (int32_t)CheckCov(tParam))); return; } @@ -199,7 +199,7 @@ GPUdic(2, 1) void GPUTPCTrackletConstructor::UpdateTracklet(int32_t /*nBlocks*/, const float sErr2 = tracker.Param().GetSystematicClusterErrorIFC2(x, tParam.GetY(), tParam.GetZ(), tracker.ISector() >= 18); err2Y += sErr2; err2Z += sErr2; - const float kFactor = tracker.Param().rec.tpc.hitPickUpFactor * tracker.Param().rec.tpc.hitPickUpFactor * 3.5f * 3.5f; + const float kFactor = tracker.GetChiSeedFactor(); float sy2 = kFactor * (tParam.Err2Y() + err2Y); float sz2 = kFactor * (tParam.Err2Z() + err2Z); if (sy2 > tracker.Param().rec.tpc.hitSearchArea2) { @@ -307,7 +307,7 @@ GPUdic(2, 1) void GPUTPCTrackletConstructor::UpdateTracklet(int32_t /*nBlocks*/, const float kFactor = tracker.Param().rec.tpc.hitPickUpFactor * tracker.Param().rec.tpc.hitPickUpFactor * 7.0f * 7.0f; const float maxWindow2 = tracker.Param().rec.tpc.hitSearchArea2; const float sy2 = CAMath::Min(maxWindow2, kFactor * (tParam.Err2Y() + err2Y)); - const float sz2 = CAMath::Min(maxWindow2, kFactor * (tParam.Err2Z() + err2Z)); + const float sz2 = CAMath::Min(maxWindow2, kFactor * (tParam.Err2Z() + err2Z)); // TODO: Double check if the factors are enough, errors could be larger with flags int32_t bin, ny, nz; row.Grid().GetBinArea(yUncorrected, zUncorrected + tParam.ZOffset(), CAMath::Sqrt(sy2), CAMath::Sqrt(sz2), bin, ny, nz); diff --git a/GPU/GPUTracking/SectorTracker/GPUTPCTrackletSelector.cxx b/GPU/GPUTracking/SectorTracker/GPUTPCTrackletSelector.cxx index 1f85718de6fa5..0a68a64c9a204 100644 --- a/GPU/GPUTracking/SectorTracker/GPUTPCTrackletSelector.cxx +++ b/GPU/GPUTracking/SectorTracker/GPUTPCTrackletSelector.cxx @@ -51,7 +51,7 @@ GPUdii() void GPUTPCTrackletSelector::Thread<0>(int32_t nBlocks, int32_t nThread uint32_t gap = 0; uint32_t nShared = 0; uint32_t nHits = 0; - const uint32_t minHits = tracker.Param().rec.tpc.minNClustersTrackSeed == -1 ? tracker.Param().tpcMinHitsB5(tracklet.Param().QPt() * tracker.Param().qptB5Scaler) : tracker.Param().rec.tpc.minNClustersTrackSeed; + const uint32_t minHits = tracker.Param().rec.tpc.minNClustersTrackSeed == -1 ? tracker.Param().tpcMinHitsB5Abs(CAMath::Abs(tracklet.Param().QPt() * tracker.Param().qptB5Scaler)) : tracker.Param().rec.tpc.minNClustersTrackSeed; const uint32_t sharingMinNorm = minHits * tracker.Param().rec.tpc.trackletMinSharedNormFactor; const float maxSharedNorm = maxSharedFrac * sharingMinNorm; diff --git a/GPU/GPUTracking/Standalone/Benchmark/CMakeLists.txt b/GPU/GPUTracking/Standalone/Benchmark/CMakeLists.txt index d3bef7c0be1f6..efd165047fb7f 100644 --- a/GPU/GPUTracking/Standalone/Benchmark/CMakeLists.txt +++ b/GPU/GPUTracking/Standalone/Benchmark/CMakeLists.txt @@ -32,6 +32,7 @@ if(ALIGPU_BUILD_TYPE STREQUAL "Standalone") endif() install(DIRECTORY ../tools DESTINATION ${CMAKE_INSTALL_DATAROOTDIR}/GPU) +install(DIRECTORY ../scripts DESTINATION ${CMAKE_INSTALL_DATAROOTDIR}/GPU USE_SOURCE_PERMISSIONS) install(DIRECTORY ../../Definitions/Parameters/ DESTINATION ${CMAKE_INSTALL_DATAROOTDIR}/GPU/tools FILES_MATCHING REGEX "\\.(python|sh|cmake)") target_compile_definitions(${targetName} PRIVATE $) diff --git a/GPU/GPUTracking/Standalone/Benchmark/standalone.cxx b/GPU/GPUTracking/Standalone/Benchmark/standalone.cxx index e831163ac7204..dff0f9462d485 100644 --- a/GPU/GPUTracking/Standalone/Benchmark/standalone.cxx +++ b/GPU/GPUTracking/Standalone/Benchmark/standalone.cxx @@ -524,7 +524,7 @@ int32_t SetupReconstruction() if (recSet.tpc.rejectionStrategy >= GPUSettings::RejectionStrategyB) { procSet.tpcInputWithClusterRejection = 1; } - recSet.tpc.disableRefitAttachment = 0xFF; + recSet.tpc.disableMarkAdjacent = 0xFF; recSet.maxTrackQPtB5 = CAMath::Min(recSet.maxTrackQPtB5, recSet.tpc.rejectQPtB5); GPUChainTracking::ApplySyncSettings(procSet, recSet, steps.steps, false, configStandalone.rundEdx); recAsync->SetSettings(&grp, &recSet, &procSet, &steps); @@ -982,7 +982,7 @@ int32_t main(int argc, char** argv) nEventsProcessed++; if (configStandalone.timeFrameTime) { - double nClusters = chainTracking->GetProcessors()->tpcMerger.NMaxClusters(); + double nClusters = chainTracking->GetProcessors()->tpcMerger.NClusters(); if (nClusters > 0) { const int32_t nOrbits = 32; const double colRate = 50000; diff --git a/GPU/GPUTracking/Standalone/scripts/compareTrackNcl.py b/GPU/GPUTracking/Standalone/scripts/compareTrackNcl.py new file mode 100755 index 0000000000000..03d863dbb55ef --- /dev/null +++ b/GPU/GPUTracking/Standalone/scripts/compareTrackNcl.py @@ -0,0 +1,40 @@ +#!/usr/bin/env python3 +# Copyright 2019-2026 CERN and copyright holders of ALICE O2. +# See https://alice-o2.web.cern.ch/copyright for details of the copyright holders. +# All rights not expressly granted are reserved. +# +# This software is distributed under the terms of the GNU General Public +# License v3 (GPL Version 3), copied verbatim in the file "COPYING". +# +# 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. + +# Usage: compareTrackNcl.py FILE1 FILE2, with files of lines "Track NCl " (from --DEBUGmergerDumpTrackNCl). +# Prints the ids of tracks with less than half the clusters in FILE1 compared to FILE2 as DEBUG_TRACK="...", to be used as environment variable with GPUCA_CADEBUG 2 in GPUTPCGMTrackParam.cxx, followed by some statistics. + +import sys + + +def readNCl(fileName): + ncl = {} + with open(fileName) as f: + for line in f: + cols = line.split() + if len(cols) >= 4 and cols[0] == "Track" and cols[2] == "NCl": + ncl[int(cols[1])] = int(cols[3]) + return ncl + + +if len(sys.argv) != 3: + sys.exit("Usage: " + sys.argv[0] + " FILE1 FILE2") +ncl1, ncl2 = readNCl(sys.argv[1]), readNCl(sys.argv[2]) +common = [i for i in sorted(ncl1) if i in ncl2] +ids = [str(i) for i in common if 2 * ncl1[i] < ncl2[i]] +print('DEBUG_TRACK="' + "\n".join(",".join(ids[i:i + 30]) for i in range(0, len(ids), 30)) + '"') +print() # Statistics as shell comments, so that the output can still be passed to eval +print("# Tracks in total (present in both files):", len(common)) +for text, n in (("Tracks with 0 clusters in file 2", sum(1 for i in common if ncl2[i] == 0)), + ("Tracks with 0 clusters in file 1 and > 0 in file 2", sum(1 for i in common if ncl1[i] == 0 and ncl2[i] > 0)), + ("Tracks with NCl(file 1) < 0.5 * NCl(file 2)", len(ids))): + print("# %s: %d (%.2f%%)" % (text, n, 100. * n / max(1, len(common)))) diff --git a/GPU/GPUTracking/TRDTracking/GPUTRDInterfaces.h b/GPU/GPUTracking/TRDTracking/GPUTRDInterfaces.h index 0d71b35e92c78..6ef8643215e6d 100644 --- a/GPU/GPUTracking/TRDTracking/GPUTRDInterfaces.h +++ b/GPU/GPUTracking/TRDTracking/GPUTRDInterfaces.h @@ -147,7 +147,7 @@ class trackInterface : public GPUTPCGMTrackParam GPUd() float getSnp() const { return GetSinPhi(); } GPUd() float getTgl() const { return GetDzDs(); } GPUd() float getQ2Pt() const { return GetQPt(); } - GPUd() float getEta() const { return -CAMath::Log(CAMath::Tan(0.5f * (0.5f * M_PI - CAMath::ATan(getTgl())))); } + GPUd() float getEta() const { return -CAMath::Log(CAMath::Tan(0.5f * (0.5f * CAMath::Pi() - CAMath::ATan(getTgl())))); } GPUd() float getPt() const { return CAMath::Abs(getQ2Pt()) > 0 ? CAMath::Abs(1.f / getQ2Pt()) : 99999.f; } GPUd() float getSigmaY2() const { return GetErr2Y(); } GPUd() float getSigmaZ2() const { return GetErr2Z(); } diff --git a/GPU/GPUTracking/kernels.cmake b/GPU/GPUTracking/kernels.cmake index ed155788b0bef..539840baaff7a 100644 --- a/GPU/GPUTracking/kernels.cmake +++ b/GPU/GPUTracking/kernels.cmake @@ -51,7 +51,7 @@ o2_gpu_add_kernel("GPUTPCGlobalDebugSortKernels, mergedTracks2" "= TPC o2_gpu_add_kernel("GPUTPCGlobalDebugSortKernels, borderTracks" "= TPCMERGER" NO int8_t parameter) o2_gpu_add_kernel("GPUTPCCreateOccupancyMap, fill" "= TPCOCCUPANCY" LB GPUTPCClusterOccupancyMapBin* map) o2_gpu_add_kernel("GPUTPCCreateOccupancyMap, fold" "= TPCOCCUPANCY" LB GPUTPCClusterOccupancyMapBin* map uint32_t* output) -o2_gpu_add_kernel("GPUTPCGMMergerTrackFit" "GPUTPCGMMergerGPU TPCMERGER TPCTRACKER MATLUT TPCDEDX" LB int32_t mode) +o2_gpu_add_kernel("GPUTPCGMMergerTrackFit" "GPUTPCGMMergerGPU TPCMERGER TPCTRACKER MATLUT TPCDEDX" LB int32_t mode int32_t rebuilt) o2_gpu_add_kernel("GPUTPCGMMergerFollowLoopers" "GPUTPCGMMergerGPU TPCMERGER TPCTRACKER MATLUT" LB) o2_gpu_add_kernel("GPUTPCGMMergerUnpackResetIds" "GPUTPCGMMergerGPU TPCMERGER" LB int32_t iSector) o2_gpu_add_kernel("GPUTPCGMMergerSectorRefit" "GPUTPCGMMergerGPU TPCMERGER MATLUT" LB int32_t iSector) @@ -84,6 +84,11 @@ o2_gpu_add_kernel("GPUTPCGMMergerFinalize, step2" "GPUTP o2_gpu_add_kernel("GPUTPCGMMergerMergeLoopers, step0" "GPUTPCGMMergerGPU TPCMERGER" LB) o2_gpu_add_kernel("GPUTPCGMMergerMergeLoopers, step1" "GPUTPCGMMergerGPU TPCMERGER" LB) o2_gpu_add_kernel("GPUTPCGMMergerMergeLoopers, step2" "GPUTPCGMMergerGPU TPCMERGER" LB) +o2_gpu_add_kernel("GPUTPCGMMergerHitWeights, prepare" "GPUTPCGMMergerGPU TPCMERGER" LB int32_t iteration) +o2_gpu_add_kernel("GPUTPCGMMergerHitWeights, compute" "GPUTPCGMMergerGPU TPCMERGER" LB int32_t iteration) +o2_gpu_add_kernel("GPUTPCGMMergerHitWeights, resolve1" "GPUTPCGMMergerGPU TPCMERGER" LB int32_t iteration) +o2_gpu_add_kernel("GPUTPCGMMergerHitWeights, resolve2" "GPUTPCGMMergerGPU TPCMERGER" LB int32_t iteration) +o2_gpu_add_kernel("GPUTPCGMMergerHitWeights, resolveShared" "GPUTPCGMMergerGPU TPCMERGER" LB int32_t iteration) o2_gpu_add_kernel("GPUTPCGMO2Output, prepare" "= TPCMERGER" LB) o2_gpu_add_kernel("GPUTPCGMO2Output, sort" "= TPCMERGER") o2_gpu_add_kernel("GPUTPCGMO2Output, output" "= TPCMERGER" LB) diff --git a/GPU/GPUTracking/qa/GPUQA.cxx b/GPU/GPUTracking/qa/GPUQA.cxx index 86aaea47b57f0..2038c4c84f707 100644 --- a/GPU/GPUTracking/qa/GPUQA.cxx +++ b/GPU/GPUTracking/qa/GPUQA.cxx @@ -182,7 +182,7 @@ static const constexpr int32_t COLORS_HEX[COLORCOUNT] = {0xB03030, 0x00A000, 0x0 static const constexpr int32_t CONFIG_DASHED_MARKERS = 0; static const constexpr float AXES_MIN[5] = {-Y_MAX, -Z_MAX, 0.f, -ETA_MAX, PT_MIN}; -static const constexpr float AXES_MAX[5] = {Y_MAX, Z_MAX, 2.f * M_PI, ETA_MAX, PT_MAX}; +static const constexpr float AXES_MAX[5] = {Y_MAX, Z_MAX, 2.f * CAMath::Pi(), ETA_MAX, PT_MAX}; static const constexpr int32_t AXIS_BINS[5] = {51, 51, 144, 31, 50}; static const constexpr int32_t RES_AXIS_BINS[] = {1017, 113}; // Consecutive bin sizes, histograms are binned down until the maximum entry is 50, each bin size should evenly divide its predecessor. static const constexpr float RES_AXES[5] = {1., 1., 0.03, 0.03, 1.0}; @@ -1247,12 +1247,12 @@ void GPUQA::RunQA(bool matchOnly, const std::vector* tracksEx const mcInfo_t& info = GetMCTrack(i, iCol); additionalMCParameters& mc2 = mMCParam[iCol][i]; mc2.pt = std::sqrt(info.pX * info.pX + info.pY * info.pY); - mc2.phi = M_PI + std::atan2(-info.pY, -info.pX); + mc2.phi = CAMath::Pi() + std::atan2(-info.pY, -info.pX); float p = info.pX * info.pX + info.pY * info.pY + info.pZ * info.pZ; if (p < 1e-18) { mc2.theta = mc2.eta = 0.f; } else { - mc2.theta = info.pZ == 0 ? (M_PI / 2) : (std::acos(info.pZ / std::sqrt(p))); + mc2.theta = info.pZ == 0 ? (CAMath::Pi() / 2) : (std::acos(info.pZ / std::sqrt(p))); mc2.eta = -std::log(std::tan(0.5 * mc2.theta)); } if (mConfig.writeMCLabels) { @@ -1306,10 +1306,10 @@ void GPUQA::RunQA(bool matchOnly, const std::vector* tracksEx } float alpha = std::atan2(info.y, info.x); - alpha /= M_PI / 9.f; + alpha /= CAMath::Pi() / 9.f; alpha = std::floor(alpha); - alpha *= M_PI / 9.f; - alpha += M_PI / 18.f; + alpha *= CAMath::Pi() / 9.f; + alpha += CAMath::Pi() / 18.f; float c = std::cos(alpha); float s = std::sin(alpha); @@ -1870,7 +1870,7 @@ void GPUQA::RunQA(bool matchOnly, const std::vector* tracksEx } } - uint32_t nCl = clNative ? clNative->nClustersTotal : mTracking->GetProcessors()->tpcMerger.NMaxClusters(); + uint32_t nCl = clNative ? clNative->nClustersTotal : mTracking->GetProcessors()->tpcMerger.NClusters(); mClusterCounts.nTotal += nCl; if (mQATasks & (taskClusterCounts | taskClusterRejection)) { for (uint32_t iSector = 0; iSector < GPUTPCGeometry::NSECTORS; iSector++) { @@ -2968,7 +2968,7 @@ int32_t GPUQA::DrawQAHistograms(TObjArray* qcout) } } - if (mQATasks & taskTrackStatistics) { + if (mQATasks & taskTrackStatistics) { // TODO: Add plot for Chi2/NDF, and NCl correct and fake MC histograms // Process track statistic histograms float tmpMax = 0.; diff --git a/GPU/GPUTracking/qa/GPUQAHelper.h b/GPU/GPUTracking/qa/GPUQAHelper.h index dcbe7a94711aa..6712fd0a603ed 100644 --- a/GPU/GPUTracking/qa/GPUQAHelper.h +++ b/GPU/GPUTracking/qa/GPUQAHelper.h @@ -162,6 +162,17 @@ static inline auto GPUTPCTrkLbl(const AliHLTTPCClusterMCLabel* x, Args... args) } } +template +static inline bool GPUTPCTrkLblSearch(const T& clusterLabels, const MCCompLabel& trkLabel) +{ + for (const auto& clLabel : clusterLabels) { + if (trkLabel.compare(clLabel) >= 0) { + return true; + } + } + return false; +} + } // namespace gpu } // namespace o2 diff --git a/GPU/GPUTracking/qa/genEvents.cxx b/GPU/GPUTracking/qa/genEvents.cxx index 38c2cf2914770..c1f33a1e0bbc9 100644 --- a/GPU/GPUTracking/qa/genEvents.cxx +++ b/GPU/GPUTracking/qa/genEvents.cxx @@ -202,8 +202,8 @@ int32_t genEvents::GenerateEvent(const GPUParam& param, const char* filename) double eta = gRandom->Uniform(-1.5, 1.5); double theta = 2 * std::atan(1. / exp(eta)); - double lambda = theta - M_PI / 2; - // double theta = gRandom->Uniform(-60,60)*M_PI/180.; + double lambda = theta - CAMath::Pi() / 2; + // double theta = gRandom->Uniform(-60,60)*CAMath::Pi()/180.; double pt = .08 * std::pow(10, gRandom->Uniform(0, 2.2)); double q = 1.; diff --git a/GPU/GPUTracking/qa/genEvents.h b/GPU/GPUTracking/qa/genEvents.h index ea209a0a46cb8..6bd6d4b370f64 100644 --- a/GPU/GPUTracking/qa/genEvents.h +++ b/GPU/GPUTracking/qa/genEvents.h @@ -16,7 +16,7 @@ #define GENEVENTS_H #include "GPUCommonDef.h" -#include +#include "GPUCommonMath.h" namespace o2::gpu { @@ -66,7 +66,7 @@ class genEvents uint32_t id; }; - const double mTwoPi = 2 * M_PI; + const double mTwoPi = 2 * CAMath::Pi(); const double mSectorDAngle = mTwoPi / 18.; const double mSectorAngleOffset = mSectorDAngle / 2;