diff --git a/Common/ML/src/OrtInterface.cxx b/Common/ML/src/OrtInterface.cxx index 58d80eb9c0bf0..d30d05d1d1a00 100644 --- a/Common/ML/src/OrtInterface.cxx +++ b/Common/ML/src/OrtInterface.cxx @@ -54,7 +54,7 @@ void OrtModel::initOptions(std::unordered_map optionsM // Load from options map if (!optionsMap.contains("model-path")) { - LOG(fatal) << "(ORT) Model path cannot be empty!"; + LOG(fatal) << "(ORT) Model path must be contained in options map!"; } if (!optionsMap["model-path"].empty()) { diff --git a/GPU/GPUTracking/Definitions/GPUSettingsList.h b/GPU/GPUTracking/Definitions/GPUSettingsList.h index 65697b7f7c08b..6419d63bb7ada 100644 --- a/GPU/GPUTracking/Definitions/GPUSettingsList.h +++ b/GPU/GPUTracking/Definitions/GPUSettingsList.h @@ -259,7 +259,7 @@ AddOption(nnInferenceEnableOrtOptimization, unsigned int, 99, "", 0, "Enables gr AddOption(nnInferenceUseDeterministicCompute, int, 0, "", 0, "Enables deterministic compute in ONNX Runtime were possible. Can be [0, 1] -> see https://github.com/microsoft/onnxruntime/blob/3b97d79b3c12dbf93aa0d563f345714596dc8ab6/onnxruntime/core/framework/session_options.h#L208") AddOption(nnInferenceOrtProfiling, int, 0, "", 0, "Enables profiling of model execution in ONNX Runtime") AddOption(nnInferenceOrtProfilingPath, std::string, ".", "", 0, "If nnInferenceOrtProfiling is set, the path to store the profiling data") -AddOption(nnInferenceVerbosity, int, 1, "", 0, "0: No messages; 1: Warnings; 2: Warnings + major debugs; >3: All debugs") +AddOption(nnInferenceVerbosity, int, 2, "", 0, "0: All debugs; 1: Warnings + major debugs; 2: Warnings; >=3: No messages") AddOption(nnClusterizerAddIndexData, int, 1, "", 0, "If normalized index data (sector, row, pad), should be appended to the input") AddOption(nnClusterizerSizeInputRow, int, 3, "", 0, "Size of the input to the NN (currently calcualted as (length-1)/2") AddOption(nnClusterizerSizeInputPad, int, 3, "", 0, "Size of the input to the NN (currently calcualted as (length-1)/2") @@ -267,7 +267,7 @@ AddOption(nnClusterizerSizeInputTime, int, 3, "", 0, "Size of the input to the N AddOption(nnClusterizerUseCfRegression, int, 0, "", 0, "(bool, default = false) If true, use the regression from the native clusterizer and not the NN") AddOption(nnClusterizerApplyCfDeconvolution, int, 0, "", 0, "Applies the CFDeconvolution kernel before the digits to the network are filled") AddOption(nnClusterizerBatchedMode, unsigned int, 1, "", 0, "(int, default = 1) If >1, the NN is evaluated on batched input of size specified in this variable") -AddOption(nnClusterizerVerbosity, int, -1, "", 0, "(int, default = -1) If >0, logging messages of the clusterizer will be displayed") +AddOption(nnClusterizerVerbosity, int, -1, "", 0, "(int, default = -1) If >0, logging messages of the clusterizer will be displayed. Higher number = higher verbosity") AddOption(nnClusterizerBoundaryFillValue, int, -1, "", 0, "Fill value for the boundary of the input to the NN") AddOption(nnClusterizerApplyNoiseSuppression, int, 1, "", 0, "Applies the NoiseSuppression kernel before the digits to the network are filled") AddOption(nnClusterizerSetDeconvolutionFlags, int, 1, "", 0, "Runs the deconvolution kernel without overwriting the charge in order to make cluster-to-track attachment identical to heuristic CF") diff --git a/GPU/GPUTracking/Global/GPUChainTrackingClusterizer.cxx b/GPU/GPUTracking/Global/GPUChainTrackingClusterizer.cxx index c92049b040c46..fd3699ae4d125 100644 --- a/GPU/GPUTracking/Global/GPUChainTrackingClusterizer.cxx +++ b/GPU/GPUTracking/Global/GPUChainTrackingClusterizer.cxx @@ -43,7 +43,6 @@ #include "DataFormatsTPC/Digit.h" #include "DataFormatsTPC/Constants.h" #include "TPCBase/RDHUtils.h" -#include "GPULogging.h" #ifdef GPUCA_HAS_ONNX #include "GPUTPCNNClusterizerKernels.h" @@ -706,7 +705,7 @@ int32_t GPUChainTracking::RunTPCClusterizer(bool synchronizeOutput) // nnApplications[lane].directOrtAllocator((nnApplications[lane].mModelClass).getEnv(), (nnApplications[lane].mModelClass).getMemoryInfo(), mRec, recreateMemoryAllocator); (nnApplications[lane].mModelReg2).initSession(); } - if (nn_settings.nnClusterizerVerbosity < 3) { + if (nn_settings.nnClusterizerVerbosity > 0) { LOG(info) << "(ORT) Allocated ONNX stream for lane " << lane << " and device " << deviceId; } }); @@ -724,12 +723,24 @@ int32_t GPUChainTracking::RunTPCClusterizer(bool synchronizeOutput) clustererNNShadow.mNnClusterizerTotalClusters = processors()->tpcClusterer[lane].mNMaxClusters; nnApplications[lane].initClusterizer(nn_settings, clustererNNShadow); } + if (nn_settings.nnClusterizerVerbosity > 2) { + LOG(info) << "(NNCLUS, GPUChainTrackingClusterizer, this=" << this << ") Processor initialized. Sector " << sector << ", lane " << lane << ", max clusters " << clustererNN.mNnClusterizerTotalClusters << " (clustererNN=" << &clustererNN << ", clustererNNShadow=" << &clustererNNShadow << ")"; + } AllocateRegisteredMemory(clustererNN.mMemoryId); + if (nn_settings.nnClusterizerVerbosity > 2) { + LOG(info) << "(NNCLUS, GPUChainTrackingClusterizer, this=" << this << ") Memory registered for memoryId " << clustererNN.mMemoryId << " (clustererNN=" << &clustererNN << ", clustererNNShadow=" << &clustererNNShadow << ")"; + } // nnApplications[lane].createBoundary(clustererNNShadow); // nnApplications[lane].createIndexLookup(clustererNNShadow); } if (doGPU) { + if (nn_settings.nnClusterizerVerbosity > 2) { + LOG(info) << "(NNCLUS, GPUChainTrackingClusterizer, this=" << this << ") Writing to constant memory..."; + } WriteToConstantMemory(RecoStep::TPCClusterFinding, (char*)&processors()->tpcNNClusterer - (char*)processors(), &processorsShadow()->tpcNNClusterer, sizeof(GPUTPCNNClusterizer) * NSECTORS, mRec->NStreams() - 1, &mEvents->init); + if (nn_settings.nnClusterizerVerbosity > 2) { + LOG(info) << "(NNCLUS, GPUChainTrackingClusterizer, this=" << this << ") Writing to constant memory done"; + } } } #endif @@ -1010,7 +1021,13 @@ int32_t GPUChainTracking::RunTPCClusterizer(bool synchronizeOutput) } // float time_clusterizer = 0, time_fill = 0, time_networks = 0; + if (nn_settings.nnClusterizerVerbosity > 2) { + LOG(info) << "(NNCLUS, GPUChainTrackingClusterizer, this=" << this << ") Starting loop over batched data. clustererNNShadow.mNnClusterizerBatchedMode=" << clustererNNShadow.mNnClusterizerBatchedMode << ", numLoops=" << std::ceil((float)clusterer.mPmemory->counters.nClusters / clustererNNShadow.mNnClusterizerBatchedMode) << ", numClusters=" << clusterer.mPmemory->counters.nClusters << ". (clustererNN=" << &clustererNN << ", clustererNNShadow=" << &clustererNNShadow << ")"; + } for (int batch = 0; batch < std::ceil((float)clusterer.mPmemory->counters.nClusters / clustererNNShadow.mNnClusterizerBatchedMode); batch++) { + if (nn_settings.nnClusterizerVerbosity > 3) { + LOG(info) << "(NNCLUS, GPUChainTrackingClusterizer, this=" << this << ") Start. Loop=" << batch << ". (clustererNN=" << &clustererNN << ", clustererNNShadow=" << &clustererNNShadow << ")"; + } uint batchStart = batch * clustererNNShadow.mNnClusterizerBatchedMode; size_t iSize = CAMath::Min((uint)clustererNNShadow.mNnClusterizerBatchedMode, (uint)(clusterer.mPmemory->counters.nClusters - batchStart)); @@ -1022,9 +1039,15 @@ int32_t GPUChainTracking::RunTPCClusterizer(bool synchronizeOutput) // Fills the whole input matrix at once -> better performance on CPU, but worse parallelizability runKernel({GetGrid(iSize, lane), krnlRunRangeNone}, iSector, clustererNNShadow.mNnInferenceInputDType, propagateMCLabels, batchStart); } + if (nn_settings.nnClusterizerVerbosity > 3) { + LOG(info) << "(NNCLUS, GPUChainTrackingClusterizer, this=" << this << ") Done filling data. Loop=" << batch << ". (clustererNN=" << &clustererNN << ", clustererNNShadow=" << &clustererNNShadow << ")"; + } if (clustererNNShadow.mNnClusterizerSetDeconvolutionFlags) { runKernel({GetGrid(iSize, lane), krnlRunRangeNone}, iSector, clustererNNShadow.mNnInferenceInputDType, propagateMCLabels, batchStart); // Publishing the deconvolution flags + if (nn_settings.nnClusterizerVerbosity > 3) { + LOG(info) << "(NNCLUS, GPUChainTrackingClusterizer, this=" << this << ") Done setting deconvolution flags. Loop=" << batch << ". (clustererNN=" << &clustererNN << ", clustererNNShadow=" << &clustererNNShadow << ")"; + } } // NN evaluations @@ -1044,6 +1067,9 @@ int32_t GPUChainTracking::RunTPCClusterizer(bool synchronizeOutput) } } if(GetProcessingSettings().debugLevel >= 1 && doGPU) { nnTimers[3*lane]->Stop(); } + if (nn_settings.nnClusterizerVerbosity > 3) { + LOG(info) << "(NNCLUS, GPUChainTrackingClusterizer, this=" << this << ") Done with NN classification inference. Loop=" << batch << ". (clustererNN=" << &clustererNN << ", clustererNNShadow=" << &clustererNNShadow << ")"; + } } if (!clustererNNShadow.mNnClusterizerUseCfRegression) { if(GetProcessingSettings().debugLevel >= 1 && doGPU) { nnTimers[3*lane + 1]->Start(); } @@ -1078,9 +1104,13 @@ int32_t GPUChainTracking::RunTPCClusterizer(bool synchronizeOutput) } if(GetProcessingSettings().debugLevel >= 1 && doGPU) { nnTimers[3*lane + 2]->Stop(); } } + if (nn_settings.nnClusterizerVerbosity > 3) { + LOG(info) << "(NNCLUS, GPUChainTrackingClusterizer, this=" << this << ") Done with NN regression inference. Loop=" << batch << ". (clustererNN=" << &clustererNN << ", clustererNNShadow=" << &clustererNNShadow << ")"; + } } // Publishing kernels for class labels and regression results + // In case classification should not be used, this kernel should still be executed to fill the mOutputDataClass array with default values if (nnApplication.mModelClass.getNumOutputNodes()[0][1] == 1) { runKernel({GetGrid(iSize, lane), krnlRunRangeNone}, iSector, clustererNNShadow.mNnInferenceOutputDType, propagateMCLabels, batchStart); // Assigning class labels } else { @@ -1092,6 +1122,9 @@ int32_t GPUChainTracking::RunTPCClusterizer(bool synchronizeOutput) runKernel({GetGrid(iSize, lane), krnlRunRangeNone}, iSector, clustererNNShadow.mNnInferenceOutputDType, propagateMCLabels, batchStart); // Publishing class 2 regression results } } + if (nn_settings.nnClusterizerVerbosity > 3) { + LOG(info) << "(NNCLUS, GPUChainTrackingClusterizer, this=" << this << ") Done publishing. Loop=" << batch << ". (clustererNN=" << &clustererNN << ", clustererNNShadow=" << &clustererNNShadow << ")"; + } } if (clustererNNShadow.mNnClusterizerUseCfRegression) { @@ -1100,6 +1133,9 @@ int32_t GPUChainTracking::RunTPCClusterizer(bool synchronizeOutput) } DoDebugAndDump(RecoStep::TPCClusterFinding, GPUChainTrackingDebugFlags::TPCClustererChargeMap, clusterer, &GPUTPCClusterFinder::DumpChargeMap, *mDebugFile, "Split Charges"); runKernel({GetGrid(clusterer.mPmemory->counters.nClusters, lane), krnlRunRangeNone}, iSector, clustererNNShadow.mNnInferenceInputDType, propagateMCLabels, 0); // Running the CF regression kernel - no batching needed: batchStart = 0 + if (nn_settings.nnClusterizerVerbosity > 3) { + LOG(info) << "(NNCLUS, GPUChainTrackingClusterizer, this=" << this << ") Done with CF regression. (clustererNN=" << &clustererNN << ", clustererNNShadow=" << &clustererNNShadow << ")"; + } } #else GPUFatal("Project not compiled with neural network clusterization. Aborting."); @@ -1203,7 +1239,9 @@ int32_t GPUChainTracking::RunTPCClusterizer(bool synchronizeOutput) for (int32_t i = 0; i < GetProcessingSettings().nTPCClustererLanes; i++) { #ifdef GPUCA_HAS_ONNX if (GetProcessingSettings().nn.applyNNclusterizer) { - LOG(info) << "(ORT) Environment releasing..."; + if (GetProcessingSettings().nn.nnClusterizerVerbosity > 0) { + LOG(info) << "(ORT) Environment releasing..."; + } GPUTPCNNClusterizerHost& nnApplication = nnApplications[i]; nnApplication.mModelClass.release(true); nnApplication.mModelReg1.release(true); diff --git a/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFClusterizer.cxx b/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFClusterizer.cxx index c9c6b157499f2..49ee5957b8b36 100644 --- a/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFClusterizer.cxx +++ b/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFClusterizer.cxx @@ -35,5 +35,5 @@ GPUdii() void GPUTPCCFClusterizer::Thread<0>(int32_t nBlocks, int32_t nThreads, tpc::ClusterNative* clusterOut = (onlyMC) ? nullptr : clusterer.mPclusterByRow; - GPUTPCCFClusterizer::computeClustersImpl(get_num_groups(0), get_local_size(0), get_group_id(0), get_local_id(0), clusterer, clusterer.mPmemory->fragment, smem, chargeMap, clusterer.mPfilteredPeakPositions, clusterer.Param().rec, CPU_PTR(&labelAcc), clusterer.mPmemory->counters.nClusters, clusterer.mNMaxClusterPerRow, clusterer.mPclusterInRow, clusterOut, clusterer.mPclusterPosInRow); + GPUTPCCFClusterizer::computeClustersImpl(get_num_groups(0), get_local_size(0), get_group_id(0), get_local_id(0), clusterer, clusterer.mPmemory->fragment, smem, chargeMap, clusterer.mPfilteredPeakPositions, clusterer.Param().rec, CPU_PTR(&labelAcc), clusterer.mPmemory->counters.nClusters, clusterer.mNMaxClusterPerRow, clusterer.mPclusterInRow, clusterOut, clusterer.mPclusterPosInRow, true); } diff --git a/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFClusterizer.h b/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFClusterizer.h index 466d13d3254de..70e21db81756c 100644 --- a/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFClusterizer.h +++ b/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFClusterizer.h @@ -57,7 +57,7 @@ class GPUTPCCFClusterizer : public GPUKernelTemplate template GPUd() static void Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUSharedMemory& smem, processorType& clusterer, int8_t); - static GPUd() void computeClustersImpl(int32_t, int32_t, int32_t, int32_t, processorType&, const CfFragment&, GPUSharedMemory&, const CfArray2D&, const CfChargePos*, const GPUSettingsRec&, MCLabelAccumulator*, uint32_t, uint32_t, uint32_t*, tpc::ClusterNative*, uint32_t*); + static GPUd() void computeClustersImpl(int32_t, int32_t, int32_t, int32_t, processorType&, const CfFragment&, GPUSharedMemory&, const CfArray2D&, const CfChargePos*, const GPUSettingsRec&, MCLabelAccumulator*, uint32_t, uint32_t, uint32_t*, tpc::ClusterNative*, uint32_t*, int8_t); static GPUd() void buildCluster(const GPUSettingsRec&, const CfArray2D&, CfChargePos, CfChargePos*, PackedCharge*, uint8_t*, ClusterAccumulator*, MCLabelAccumulator*); diff --git a/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFClusterizer.inc b/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFClusterizer.inc index e32abbf37584f..c2c104809990e 100644 --- a/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFClusterizer.inc +++ b/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFClusterizer.inc @@ -27,7 +27,8 @@ GPUdii() void GPUTPCCFClusterizer::computeClustersImpl(int32_t nBlocks, int32_t uint32_t maxClusterPerRow, uint32_t* clusterInRow, tpc::ClusterNative* clusterByRow, - uint32_t* clusterPosInRow) + uint32_t* clusterPosInRow, + int8_t isAccepted) { uint32_t idx = get_global_id(0); @@ -62,6 +63,9 @@ GPUdii() void GPUTPCCFClusterizer::computeClustersImpl(int32_t nBlocks, int32_t tpc::ClusterNative myCluster; pc.finalize(pos, charge, fragment.start); bool rejectCluster = !pc.toNative(pos, charge, myCluster, clusterer.Param(), chargeMap); + if (!isAccepted) { + rejectCluster = true; + } if (rejectCluster) { if (clusterPosInRow) { diff --git a/GPU/GPUTracking/TPCClusterFinder/GPUTPCNNClusterizer.cxx b/GPU/GPUTracking/TPCClusterFinder/GPUTPCNNClusterizer.cxx index 3dd8b0d621a56..6fac0e417ac26 100644 --- a/GPU/GPUTracking/TPCClusterFinder/GPUTPCNNClusterizer.cxx +++ b/GPU/GPUTracking/TPCClusterFinder/GPUTPCNNClusterizer.cxx @@ -16,6 +16,7 @@ #include "ML/3rdparty/GPUORTFloat16.h" #include "GPUTPCNNClusterizer.h" #include "GPUSettings.h" +#include "GPUCommonLogger.h" using namespace o2::gpu; @@ -25,6 +26,8 @@ void GPUTPCNNClusterizer::SetMaxData(const GPUTrackingInOutPointers& io) {} void* GPUTPCNNClusterizer::setIOPointers(void* mem) { + // Keep track of the start address to compute how much memory we assign + void* startMem = mem; if (mNnClusterizerBatchedMode > 0) { if (mNnInferenceInputDType == 0 && mNnClusterizerElementSize > 0) { computePointerWithAlignment(mem, mInputData_16, mNnClusterizerBatchedMode * mNnClusterizerElementSize); @@ -62,6 +65,86 @@ void* GPUTPCNNClusterizer::setIOPointers(void* mem) if (mNnClusterizerTotalClusters > 0) { computePointerWithAlignment(mem, mOutputDataClass, mNnClusterizerTotalClusters); } + + if (mNnClusterizerVerbosity > 2) { + if (mNnClusterizerVerbosity > 3) { + auto fmt = [](size_t bytes) -> const char* { + static char buf[64]; + double mb = (double)bytes / (1024.0 * 1024.0); + int n = snprintf(buf, sizeof(buf), "%zu bytes (%.3f MB)", bytes, mb); + (void)n; + return buf; + }; + + // Element counts (number of array entries, not bytes) + size_t elemsClusterFlags = (mClusterFlags && mNnClusterizerBatchedMode > 0) ? (size_t)2 * mNnClusterizerBatchedMode : 0; + size_t elemsInput16 = (mInputData_16 && mNnClusterizerBatchedMode > 0 && mNnClusterizerElementSize > 0) ? (size_t)mNnClusterizerBatchedMode * mNnClusterizerElementSize : 0; + size_t elemsInput32 = (mInputData_32 && mNnClusterizerBatchedMode > 0 && mNnClusterizerElementSize > 0) ? (size_t)mNnClusterizerBatchedMode * mNnClusterizerElementSize : 0; + size_t elemsProb16 = (mModelProbabilities_16 && mNnClusterizerBatchedMode > 0 && mNnClusterizerModelClassNumOutputNodes > 0) ? (size_t)mNnClusterizerBatchedMode * mNnClusterizerModelClassNumOutputNodes : 0; + size_t elemsProb32 = (mModelProbabilities_32 && mNnClusterizerBatchedMode > 0 && mNnClusterizerModelClassNumOutputNodes > 0) ? (size_t)mNnClusterizerBatchedMode * mNnClusterizerModelClassNumOutputNodes : 0; + size_t elemsReg1_16 = (mOutputDataReg1_16 && mNnClusterizerBatchedMode > 0 && mNnClusterizerModelReg1NumOutputNodes > 0) ? (size_t)mNnClusterizerBatchedMode * mNnClusterizerModelReg1NumOutputNodes : 0; + size_t elemsReg2_16 = (mOutputDataReg2_16 && mNnClusterizerBatchedMode > 0 && mNnClusterizerModelReg2NumOutputNodes > 0) ? (size_t)mNnClusterizerBatchedMode * mNnClusterizerModelReg2NumOutputNodes : 0; + size_t elemsReg1_32 = (mOutputDataReg1_32 && mNnClusterizerBatchedMode > 0 && mNnClusterizerModelReg1NumOutputNodes > 0) ? (size_t)mNnClusterizerBatchedMode * mNnClusterizerModelReg1NumOutputNodes : 0; + size_t elemsReg2_32 = (mOutputDataReg2_32 && mNnClusterizerBatchedMode > 0 && mNnClusterizerModelReg2NumOutputNodes > 0) ? (size_t)mNnClusterizerBatchedMode * mNnClusterizerModelReg2NumOutputNodes : 0; + size_t elemsOutputDataClass = (mOutputDataClass && mNnClusterizerTotalClusters > 0) ? (size_t)mNnClusterizerTotalClusters : 0; + + // Byte sizes + size_t szClusterFlags = elemsClusterFlags * sizeof(int8_t); + size_t szInput16 = elemsInput16 * sizeof(OrtDataType::Float16_t); + size_t szInput32 = elemsInput32 * sizeof(float); + size_t szProb16 = elemsProb16 * sizeof(OrtDataType::Float16_t); + size_t szProb32 = elemsProb32 * sizeof(float); + size_t szReg1_16 = elemsReg1_16 * sizeof(OrtDataType::Float16_t); + size_t szReg2_16 = elemsReg2_16 * sizeof(OrtDataType::Float16_t); + size_t szReg1_32 = elemsReg1_32 * sizeof(float); + size_t szReg2_32 = elemsReg2_32 * sizeof(float); + size_t szOutputDataClass = elemsOutputDataClass * sizeof(int32_t); + + LOG(info) << "(NNCLUS, GPUTPCNNClusterizer, this=" << this << ") Pointers set for clusterizer with memoryID " << mMemoryId << " deviceID " << mDeviceId << " and sector " << mISector; + LOG(info) << "(NNCLUS, GPUTPCNNClusterizer, this=" << this << ") mOutputDataClass pointer: " << mOutputDataClass + << " | elements=" << elemsOutputDataClass << " (= mNnClusterizerTotalClusters)" + << " | " << fmt(szOutputDataClass); + LOG(info) << "(NNCLUS, GPUTPCNNClusterizer, this=" << this << ") mClusterFlags pointer: " << static_cast(mClusterFlags) + << " | elements=" << elemsClusterFlags << " (= 2 * mNnClusterizerBatchedMode)" + << " | " << fmt(szClusterFlags); + LOG(info) << "(NNCLUS, GPUTPCNNClusterizer, this=" << this << ") mInputData_16 pointer: " << mInputData_16 + << " | elements=" << elemsInput16 << " (= mNnClusterizerBatchedMode * mNnClusterizerElementSize)" + << " | " << fmt(szInput16); + LOG(info) << "(NNCLUS, GPUTPCNNClusterizer, this=" << this << ") mModelProbabilities_16 pointer: " << mModelProbabilities_16 + << " | elements=" << elemsProb16 << " (= mNnClusterizerBatchedMode * mNnClusterizerModelClassNumOutputNodes)" + << " | " << fmt(szProb16); + LOG(info) << "(NNCLUS, GPUTPCNNClusterizer, this=" << this << ") mOutputDataReg1_16 pointer: " << mOutputDataReg1_16 + << " | elements=" << elemsReg1_16 << " (= mNnClusterizerBatchedMode * mNnClusterizerModelReg1NumOutputNodes)" + << " | " << fmt(szReg1_16); + LOG(info) << "(NNCLUS, GPUTPCNNClusterizer, this=" << this << ") mOutputDataReg2_16 pointer: " << mOutputDataReg2_16 + << " | elements=" << elemsReg2_16 << " (= mNnClusterizerBatchedMode * mNnClusterizerModelReg2NumOutputNodes)" + << " | " << fmt(szReg2_16); + LOG(info) << "(NNCLUS, GPUTPCNNClusterizer, this=" << this << ") mInputData_32 pointer: " << mInputData_32 + << " | elements=" << elemsInput32 << " (= mNnClusterizerBatchedMode * mNnClusterizerElementSize)" + << " | " << fmt(szInput32); + LOG(info) << "(NNCLUS, GPUTPCNNClusterizer, this=" << this << ") mModelProbabilities_32 pointer: " << mModelProbabilities_32 + << " | elements=" << elemsProb32 << " (= mNnClusterizerBatchedMode * mNnClusterizerModelClassNumOutputNodes)" + << " | " << fmt(szProb32); + LOG(info) << "(NNCLUS, GPUTPCNNClusterizer, this=" << this << ") mOutputDataReg1_32 pointer: " << mOutputDataReg1_32 + << " | elements=" << elemsReg1_32 << " (= mNnClusterizerBatchedMode * mNnClusterizerModelReg1NumOutputNodes)" + << " | " << fmt(szReg1_32); + LOG(info) << "(NNCLUS, GPUTPCNNClusterizer, this=" << this << ") mOutputDataReg2_32 pointer: " << mOutputDataReg2_32 + << " | elements=" << elemsReg2_32 << " (= mNnClusterizerBatchedMode * mNnClusterizerModelReg2NumOutputNodes)" + << " | " << fmt(szReg2_32); + } + // Compute allocated bytes (difference between advanced pointer and start pointer) + size_t allocatedBytes = static_cast(reinterpret_cast(mem) - reinterpret_cast(startMem)); + double allocatedMB = static_cast(allocatedBytes) / (1024.0 * 1024.0); + { + char allocMsg[256]; + int nn = snprintf(allocMsg, sizeof(allocMsg), + "(NNCLUS, GPUTPCNNClusterizer, this=%p) Total scratch allocation in setIOPointers: %zu bytes (%.3f MB)", + (void*)this, (size_t)allocatedBytes, allocatedMB); + (void)nn; + LOG(info) << allocMsg; + } + } + return mem; } diff --git a/GPU/GPUTracking/TPCClusterFinder/GPUTPCNNClusterizer.h b/GPU/GPUTracking/TPCClusterFinder/GPUTPCNNClusterizer.h index 69972c8a0651c..a6b0b081fc3dd 100644 --- a/GPU/GPUTracking/TPCClusterFinder/GPUTPCNNClusterizer.h +++ b/GPU/GPUTracking/TPCClusterFinder/GPUTPCNNClusterizer.h @@ -50,7 +50,7 @@ class GPUTPCNNClusterizer : public GPUProcessor int32_t mNnClusterizerUseCfRegression = 0; int32_t mNnClusterizerBatchedMode = 1; int32_t mNnClusterizerTotalClusters = 1; - int32_t mNnClusterizerVerbosity = 0; + int32_t mNnClusterizerVerbosity = 1; int32_t mNnClusterizerBoundaryFillValue = -1; int32_t mNnClusterizerModelClassNumOutputNodes = -1; int32_t mNnClusterizerModelReg1NumOutputNodes = -1; diff --git a/GPU/GPUTracking/TPCClusterFinder/GPUTPCNNClusterizerKernels.cxx b/GPU/GPUTracking/TPCClusterFinder/GPUTPCNNClusterizerKernels.cxx index 1882acd2a45c6..15ee6b6119022 100644 --- a/GPU/GPUTracking/TPCClusterFinder/GPUTPCNNClusterizerKernels.cxx +++ b/GPU/GPUTracking/TPCClusterFinder/GPUTPCNNClusterizerKernels.cxx @@ -12,7 +12,11 @@ /// \file GPUTPCNNClusterizerKernels.cxx /// \author Christian Sonnabend +#include "clusterFinderDefs.h" +#include "PackedCharge.h" #include "GPUTPCNNClusterizerKernels.h" +#include "GPUConstantMem.h" +#include "GPUTPCClusterFinder.h" #include "GPUTPCCFClusterizer.h" #include "GPUTPCGeometry.h" @@ -40,14 +44,11 @@ GPUdii() void GPUTPCNNClusterizerKernels::Thread chargeMap(reinterpret_cast(clusterer.mPchargeMap)); CPU_ONLY(MCLabelAccumulator labelAcc(clusterer)); tpc::ClusterNative* clusterOut = (withMC) ? nullptr : clusterer.mPclusterByRow; - o2::gpu::GPUTPCCFClusterizer::GPUSharedMemory smem_new; - GPUTPCCFClusterizer::computeClustersImpl(get_num_groups(0), get_local_size(0), get_group_id(0), get_local_id(0), clusterer, clusterer.mPmemory->fragment, smem_new, chargeMap, clusterer.mPfilteredPeakPositions, clusterer.Param().rec, CPU_PTR(&labelAcc), clusterer.mPmemory->counters.nClusters, clusterer.mNMaxClusterPerRow, clusterer.mPclusterInRow, clusterOut, clusterer.mPclusterPosInRow); + int8_t isAccepted = (clustererNN.mNnClusterizerUseClassification ? (clustererNN.mOutputDataClass[CAMath::Min(glo_idx, (uint32_t)clusterer.mPmemory->counters.nClusters - 1)] > 0) : 1); + GPUTPCCFClusterizer::computeClustersImpl(get_num_groups(0), get_local_size(0), get_group_id(0), get_local_id(0), clusterer, clusterer.mPmemory->fragment, reinterpret_cast(smem), chargeMap, clusterer.mPfilteredPeakPositions, clusterer.Param().rec, CPU_PTR(&labelAcc), clusterer.mPmemory->counters.nClusters, clusterer.mNMaxClusterPerRow, clusterer.mPclusterInRow, clusterOut, clusterer.mPclusterPosInRow, isAccepted); } template <> @@ -57,7 +58,7 @@ GPUdii() void GPUTPCNNClusterizerKernels::Thread= clusterer.mPmemory->counters.nClusters) { + if (glo_idx + batchStart >= clusterer.mPmemory->counters.nClusters || glo_idx >= (uint32_t)clustererNN.mNnClusterizerBatchedMode) { return; } @@ -95,12 +96,6 @@ GPUdii() void GPUTPCNNClusterizerKernels::Thread(chargeMap[tmp_pos].unpack()) / central_charge; - - if (!clustererNN.mNnClusterizerSetDeconvolutionFlags && r == 0 && CAMath::Abs(p) < 3 && CAMath::Abs(t) < 3 && p != 0 && t != 0) { - clustererNN.mClusterFlags[2 * glo_idx] += CfUtils::isPeak(isPeakMap[tmp_pos]); - clustererNN.mClusterFlags[2 * glo_idx + 1] = clustererNN.mClusterFlags[2 * glo_idx]; - } - if (dtype == 0) { clustererNN.mInputData_16[write_idx] = (OrtDataType::Float16_t)normalized_charge; } else { @@ -145,11 +140,13 @@ template <> GPUdii() void GPUTPCNNClusterizerKernels::Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUSharedMemory& smem, processorType& processors, uint8_t sector, int8_t dtype, int8_t withMC, uint32_t batchStart) { uint32_t glo_idx = get_global_id(0); - auto& clusterer = processors.tpcClusterer[sector]; auto& clustererNN = processors.tpcNNClusterer[sector]; - // Optimized division using bit operations + if (glo_idx >= (uint32_t)clustererNN.mNnClusterizerBatchedMode * clustererNN.mNnClusterizerRowTimeSizeFull) { + return; + } + uint32_t base_idx = glo_idx / clustererNN.mNnClusterizerRowTimeSizeFull; uint32_t transient_index = glo_idx - (base_idx * clustererNN.mNnClusterizerRowTimeSizeFull); @@ -185,7 +182,7 @@ GPUdii() void GPUTPCNNClusterizerKernels::Thread= clustererNN.mNnClusterizerElementSize * clustererNN.mNnClusterizerBatchedMode) { + // printf("Error: Write index out of bounds (central array)! %d >= %d (write_idx: %d, base_idx: %d, transient_index: %d, row_idx: %d, time_idx: %d, r_local: %d, t_local: %d)\n", + // write_idx, (int)(clustererNN.mNnClusterizerElementSize * clustererNN.mNnClusterizerBatchedMode), write_idx, base_idx, transient_index, row_idx, time_idx, r_local, t_local); + // } + // if ((clusterer.mPmemory->counters.nClusters - batchStart) < clustererNN.mNnClusterizerBatchedMode) { + // if (write_idx >= ((clusterer.mPmemory->counters.nClusters - batchStart) * clustererNN.mNnClusterizerElementSize)) { + // printf("Error: Write index out of bounds (end of array)! %d >= %d (write_idx: %d, base_idx: %d, transient_index: %d, row_idx: %d, time_idx: %d, r_local: %d, t_local: %d)\n", + // write_idx, (int)((clusterer.mPmemory->counters.nClusters - batchStart) * clustererNN.mNnClusterizerElementSize), write_idx, base_idx, transient_index, row_idx, time_idx, r_local, t_local); + // } + // if (write_idx > ((clusterer.mPmemory->counters.nClusters - batchStart) * clustererNN.mNnClusterizerElementSize - 5)) { + // printf("Sanity check (should appear only once) %d == %d (write_idx: %d, base_idx: %d, transient_index: %d, row_idx: %d, time_idx: %d, r_local: %d, t_local: %d)\n", + // write_idx, (int)((clusterer.mPmemory->counters.nClusters - batchStart) * clustererNN.mNnClusterizerElementSize - 4), write_idx, base_idx, transient_index, row_idx, time_idx, r_local, t_local); + // } + // } + write_idx += clustererNN.mNnClusterizerFullTimeSize; // Move to next pad position } } @@ -258,41 +270,58 @@ template <> GPUdii() void GPUTPCNNClusterizerKernels::Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUSharedMemory& smem, processorType& processors, uint8_t sector, int8_t dtype, int8_t withMC, uint32_t batchStart) { uint32_t glo_idx = get_global_id(0); - if (dtype == 0) { - processors.tpcNNClusterer[sector].mOutputDataClass[glo_idx + batchStart] = (int)((processors.tpcNNClusterer[sector].mModelProbabilities_16[glo_idx]).ToFloat() > processors.tpcNNClusterer[sector].mNnClassThreshold); - } else if (dtype == 1) { - processors.tpcNNClusterer[sector].mOutputDataClass[glo_idx + batchStart] = (int)(processors.tpcNNClusterer[sector].mModelProbabilities_32[glo_idx] > processors.tpcNNClusterer[sector].mNnClassThreshold); + auto& clusterer = processors.tpcClusterer[sector]; + auto& clustererNN = processors.tpcNNClusterer[sector]; + if (glo_idx + batchStart >= clusterer.mPmemory->counters.nClusters || glo_idx >= (uint32_t)clustererNN.mNnClusterizerBatchedMode) { + return; + } + if (clustererNN.mNnClusterizerUseClassification) { + if (dtype == 0) { + clustererNN.mOutputDataClass[glo_idx + batchStart] = (int32_t)((clustererNN.mModelProbabilities_16[glo_idx]).ToFloat() > clustererNN.mNnClassThreshold); + } else if (dtype == 1) { + clustererNN.mOutputDataClass[glo_idx + batchStart] = (int32_t)(clustererNN.mModelProbabilities_32[glo_idx] > clustererNN.mNnClassThreshold); + } + } else { + clustererNN.mOutputDataClass[glo_idx + batchStart] = 1; } } template <> GPUdii() void GPUTPCNNClusterizerKernels::Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUSharedMemory& smem, processorType& processors, uint8_t sector, int8_t dtype, int8_t withMC, uint32_t batchStart) { - auto& clustererNN = processors.tpcNNClusterer[sector]; uint32_t glo_idx = get_global_id(0); - uint32_t elem_iterator = glo_idx * clustererNN.mNnClusterizerModelClassNumOutputNodes; - float current_max_prob = 0.f; // If the neural network doesn't contain the softmax as a last layer, the outputs can range in [-infty, infty] - uint32_t class_label = 0; - for (uint32_t pIdx = elem_iterator; pIdx < elem_iterator + clustererNN.mNnClusterizerModelClassNumOutputNodes; pIdx++) { - if (pIdx == elem_iterator) { - if (dtype == 0) { - current_max_prob = static_cast(clustererNN.mModelProbabilities_16[pIdx]); - } else if (dtype == 1) { - current_max_prob = clustererNN.mModelProbabilities_32[pIdx]; - } - } else { - if (dtype == 0) { - current_max_prob = CAMath::Max(current_max_prob, clustererNN.mModelProbabilities_16[pIdx].ToFloat()); - } else if (dtype == 1) { - current_max_prob = CAMath::Max(current_max_prob, clustererNN.mModelProbabilities_32[pIdx]); + auto& clusterer = processors.tpcClusterer[sector]; + auto& clustererNN = processors.tpcNNClusterer[sector]; + if (glo_idx + batchStart >= clusterer.mPmemory->counters.nClusters || glo_idx >= (uint32_t)clustererNN.mNnClusterizerBatchedMode) { + return; + } + if (clustererNN.mNnClusterizerUseClassification) { + uint32_t elem_iterator = glo_idx * clustererNN.mNnClusterizerModelClassNumOutputNodes; + float current_max_prob = 0.f; // If the neural network doesn't contain the softmax as a last layer, the outputs can range in [-infty, infty] + uint32_t class_label = 0; + for (uint32_t pIdx = elem_iterator; pIdx < elem_iterator + clustererNN.mNnClusterizerModelClassNumOutputNodes; pIdx++) { + if (pIdx == elem_iterator) { + if (dtype == 0) { + current_max_prob = static_cast(clustererNN.mModelProbabilities_16[pIdx]); + } else if (dtype == 1) { + current_max_prob = clustererNN.mModelProbabilities_32[pIdx]; + } + } else { + if (dtype == 0) { + current_max_prob = CAMath::Max(current_max_prob, clustererNN.mModelProbabilities_16[pIdx].ToFloat()); + } else if (dtype == 1) { + current_max_prob = CAMath::Max(current_max_prob, clustererNN.mModelProbabilities_32[pIdx]); + } } } - } - // uint32_t class_label = std::distance(elem_iterator, std::max_element(elem_iterator, elem_iterator + clustererNN.mNnClusterizerModelClassNumOutputNodes)); // Multiple outputs of the class network are the probabilities for each class. The highest one "wins" - clustererNN.mOutputDataClass[glo_idx + batchStart] = class_label; - if (class_label > 1) { - clustererNN.mClusterFlags[2 * glo_idx] = 1; - clustererNN.mClusterFlags[2 * glo_idx + 1] = 1; + // uint32_t class_label = std::distance(elem_iterator, std::max_element(elem_iterator, elem_iterator + clustererNN.mNnClusterizerModelClassNumOutputNodes)); // Multiple outputs of the class network are the probabilities for each class. The highest one "wins" + clustererNN.mOutputDataClass[glo_idx + batchStart] = class_label; + if (class_label > 1) { + clustererNN.mClusterFlags[2 * glo_idx] = 1; + clustererNN.mClusterFlags[2 * glo_idx + 1] = 1; + } + } else { + clustererNN.mOutputDataClass[glo_idx + batchStart] = 1; } } @@ -302,6 +331,9 @@ GPUdii() void GPUTPCNNClusterizerKernels::Thread= (uint32_t)clustererNN.mNnClusterizerBatchedMode) { + return; + } uint32_t maxClusterNum = clusterer.mPmemory->counters.nClusters; uint32_t full_glo_idx = glo_idx + batchStart; @@ -335,88 +367,84 @@ GPUdii() void GPUTPCNNClusterizerKernels::Threadcollect(peak, central_charge)); - GPUTPCCFClusterizer::buildCluster( - clusterer.Param().rec, - chargeMap, - peak, - smem.posBcast, - smem.buf, - smem.innerAboveThreshold, - &dummy_pc, - labelAcc); - } - if ((clusterer.mPmemory->fragment).isOverlap(peak.time())) { - if (clusterer.mPclusterPosInRow) { - clusterer.mPclusterPosInRow[full_glo_idx] = clusterer.mNMaxClusterPerRow; - } - return; - } - - bool notSinglePad = false, notSingleTime = false; - for (uint16_t i = 0; i < 8; i++) { - Delta2 d = cfconsts::InnerNeighbors[i]; - CfChargePos tmp_pos = peak.delta(d); - notSinglePad |= (d.x != 0) && (static_cast(chargeMap[tmp_pos].unpack()) > 0); - notSingleTime |= (d.y != 0) && (static_cast(chargeMap[tmp_pos].unpack()) > 0); + ClusterAccumulator pc; + + // Publishing logic is taken from default clusterizer + if (withMC) { + ClusterAccumulator dummy_pc; + CPU_ONLY(labelAcc->collect(peak, central_charge)); + GPUTPCCFClusterizer::buildCluster( + clusterer.Param().rec, + chargeMap, + peak, + smem.posBcast, + smem.buf, + smem.innerAboveThreshold, + &dummy_pc, + labelAcc); + } + if ((clusterer.mPmemory->fragment).isOverlap(peak.time())) { + if (clusterer.mPclusterPosInRow) { + clusterer.mPclusterPosInRow[full_glo_idx] = clusterer.mNMaxClusterPerRow; } + return; + } - if (dtype == 0) { - pc.setFull(central_charge * clustererNN.mOutputDataReg1_16[model_output_index + 4].ToFloat(), - static_cast(peak.pad()) + clustererNN.mOutputDataReg1_16[model_output_index].ToFloat(), - notSinglePad ? clustererNN.mOutputDataReg1_16[model_output_index + 2].ToFloat() : 0.f, - (clusterer.mPmemory->fragment).start + static_cast(peak.time()) + clustererNN.mOutputDataReg1_16[model_output_index + 1].ToFloat(), - notSingleTime ? clustererNN.mOutputDataReg1_16[model_output_index + 3].ToFloat() : 0.f, - clustererNN.mClusterFlags[2 * glo_idx], - clustererNN.mClusterFlags[2 * glo_idx + 1]); - } else if (dtype == 1) { - pc.setFull(central_charge * clustererNN.mOutputDataReg1_32[model_output_index + 4], - static_cast(peak.pad()) + clustererNN.mOutputDataReg1_32[model_output_index], - notSinglePad ? clustererNN.mOutputDataReg1_32[model_output_index + 2] : 0.f, - (clusterer.mPmemory->fragment).start + static_cast(peak.time()) + clustererNN.mOutputDataReg1_32[model_output_index + 1], - notSingleTime ? clustererNN.mOutputDataReg1_32[model_output_index + 3] : 0.f, - clustererNN.mClusterFlags[2 * glo_idx], - clustererNN.mClusterFlags[2 * glo_idx + 1]); - } + // For flag influence on cluster error setting: O2/GPU/GPUTracking/Base/GPUParam.inc -> UpdateClusterError2ByState + bool notSinglePad = false, notSingleTime = false; + for (uint16_t i = 0; i < 8; i++) { + Delta2 d = cfconsts::InnerNeighbors[i]; + CfChargePos tmp_pos = peak.delta(d); + notSinglePad |= (d.x != 0) && (static_cast(chargeMap[tmp_pos].unpack()) > 0); + notSingleTime |= (d.y != 0) && (static_cast(chargeMap[tmp_pos].unpack()) > 0); + } - tpc::ClusterNative myCluster; - bool rejectCluster = !pc.toNative(peak, central_charge, myCluster, clusterer.Param(), chargeMap); - if (rejectCluster) { - if (clusterer.mPclusterPosInRow) { - clusterer.mPclusterPosInRow[full_glo_idx] = clusterer.mNMaxClusterPerRow; - } - return; - } + if (dtype == 0) { + pc.setFull(central_charge * clustererNN.mOutputDataReg1_16[model_output_index + 4].ToFloat(), + static_cast(peak.pad()) + clustererNN.mOutputDataReg1_16[model_output_index].ToFloat(), + notSinglePad ? clustererNN.mOutputDataReg1_16[model_output_index + 2].ToFloat() : 0.f, + (clusterer.mPmemory->fragment).start + static_cast(peak.time()) + clustererNN.mOutputDataReg1_16[model_output_index + 1].ToFloat(), + notSingleTime ? clustererNN.mOutputDataReg1_16[model_output_index + 3].ToFloat() : 0.f, + clustererNN.mClusterFlags[2 * glo_idx], + clustererNN.mClusterFlags[2 * glo_idx + 1]); + } else if (dtype == 1) { + pc.setFull(central_charge * clustererNN.mOutputDataReg1_32[model_output_index + 4], + static_cast(peak.pad()) + clustererNN.mOutputDataReg1_32[model_output_index], + notSinglePad ? clustererNN.mOutputDataReg1_32[model_output_index + 2] : 0.f, + (clusterer.mPmemory->fragment).start + static_cast(peak.time()) + clustererNN.mOutputDataReg1_32[model_output_index + 1], + notSingleTime ? clustererNN.mOutputDataReg1_32[model_output_index + 3] : 0.f, + clustererNN.mClusterFlags[2 * glo_idx], + clustererNN.mClusterFlags[2 * glo_idx + 1]); + } - uint32_t rowIndex = 0; - if (clusterOut != nullptr) { - rowIndex = GPUTPCCFClusterizer::sortIntoBuckets( - clusterer, - myCluster, - peak.row(), - clusterer.mNMaxClusterPerRow, - clusterer.mPclusterInRow, - clusterOut); - if (clusterer.mPclusterPosInRow != nullptr) { - clusterer.mPclusterPosInRow[full_glo_idx] = rowIndex; - } - } else if (clusterer.mPclusterPosInRow) { - rowIndex = clusterer.mPclusterPosInRow[full_glo_idx]; - } - CPU_ONLY(labelAcc->commit(peak.row(), rowIndex, clusterer.mNMaxClusterPerRow)); - } else { + tpc::ClusterNative myCluster; + bool rejectCluster = !pc.toNative(peak, central_charge, myCluster, clusterer.Param(), chargeMap); + if (clustererNN.mNnClusterizerUseClassification) { + rejectCluster |= (clustererNN.mOutputDataClass[CAMath::Min(full_glo_idx, (uint32_t)clusterer.mPmemory->counters.nClusters - 1)] <= 0); + } + if (rejectCluster) { if (clusterer.mPclusterPosInRow) { clusterer.mPclusterPosInRow[full_glo_idx] = clusterer.mNMaxClusterPerRow; } return; } + + uint32_t rowIndex = 0; + if (clusterOut != nullptr) { + rowIndex = GPUTPCCFClusterizer::sortIntoBuckets( + clusterer, + myCluster, + peak.row(), + clusterer.mNMaxClusterPerRow, + clusterer.mPclusterInRow, + clusterOut); + if (clusterer.mPclusterPosInRow != nullptr) { + clusterer.mPclusterPosInRow[full_glo_idx] = rowIndex; + } + } else if (clusterer.mPclusterPosInRow) { + rowIndex = clusterer.mPclusterPosInRow[full_glo_idx]; + } + CPU_ONLY(labelAcc->commit(peak.row(), rowIndex, clusterer.mNMaxClusterPerRow)); } template <> @@ -425,6 +453,9 @@ GPUdii() void GPUTPCNNClusterizerKernels::Thread= (uint32_t)clustererNN.mNnClusterizerBatchedMode) { + return; + } uint32_t maxClusterNum = clusterer.mPmemory->counters.nClusters; CfArray2D chargeMap(reinterpret_cast(clusterer.mPchargeMap)); @@ -455,123 +486,121 @@ GPUdii() void GPUTPCNNClusterizerKernels::Thread 0) || (clustererNN.mNnClusterizerUseClassification <= 0)) { - - ClusterAccumulator pc; - - if (withMC) { - ClusterAccumulator dummy_pc; - CPU_ONLY(labelAcc->collect(peak, central_charge)); - GPUTPCCFClusterizer::buildCluster( - clusterer.Param().rec, - chargeMap, - peak, - smem.posBcast, - smem.buf, - smem.innerAboveThreshold, - &dummy_pc, - labelAcc); - } - if ((clusterer.mPmemory->fragment).isOverlap(peak.time())) { - if (clusterer.mPclusterPosInRow) { - clusterer.mPclusterPosInRow[full_glo_idx] = clusterer.mNMaxClusterPerRow; - } - return; - } - - // Cluster 1 - if (dtype == 0) { - pc.setFull(central_charge * clustererNN.mOutputDataReg2_16[model_output_index + 8].ToFloat(), - static_cast(peak.pad()) + clustererNN.mOutputDataReg2_16[model_output_index].ToFloat(), - clustererNN.mOutputDataReg2_16[model_output_index + 4].ToFloat(), - (clusterer.mPmemory->fragment).start + static_cast(peak.time()) + clustererNN.mOutputDataReg2_16[model_output_index + 2].ToFloat(), - clustererNN.mOutputDataReg2_16[model_output_index + 6].ToFloat(), - clustererNN.mClusterFlags[2 * glo_idx], - clustererNN.mClusterFlags[2 * glo_idx + 1]); - } else if (dtype == 1) { - pc.setFull(central_charge * clustererNN.mOutputDataReg2_32[model_output_index + 8], - static_cast(peak.pad()) + clustererNN.mOutputDataReg2_32[model_output_index], - clustererNN.mOutputDataReg2_32[model_output_index + 4], - (clusterer.mPmemory->fragment).start + static_cast(peak.time()) + clustererNN.mOutputDataReg2_32[model_output_index + 2], - clustererNN.mOutputDataReg2_32[model_output_index + 6], - clustererNN.mClusterFlags[2 * glo_idx], - clustererNN.mClusterFlags[2 * glo_idx + 1]); + ClusterAccumulator pc; + + if (withMC) { + ClusterAccumulator dummy_pc; + CPU_ONLY(labelAcc->collect(peak, central_charge)); + GPUTPCCFClusterizer::buildCluster( + clusterer.Param().rec, + chargeMap, + peak, + smem.posBcast, + smem.buf, + smem.innerAboveThreshold, + &dummy_pc, + labelAcc); + } + if ((clusterer.mPmemory->fragment).isOverlap(peak.time())) { + if (clusterer.mPclusterPosInRow) { + clusterer.mPclusterPosInRow[full_glo_idx] = clusterer.mNMaxClusterPerRow; } + return; + } - tpc::ClusterNative myCluster; - bool rejectCluster = !pc.toNative(peak, central_charge, myCluster, clusterer.Param(), chargeMap); - if (rejectCluster) { - if (clusterer.mPclusterPosInRow) { - clusterer.mPclusterPosInRow[full_glo_idx] = clusterer.mNMaxClusterPerRow; - } - return; - } + // Cluster 1 + if (dtype == 0) { + pc.setFull(central_charge * clustererNN.mOutputDataReg2_16[model_output_index + 8].ToFloat(), + static_cast(peak.pad()) + clustererNN.mOutputDataReg2_16[model_output_index].ToFloat(), + clustererNN.mOutputDataReg2_16[model_output_index + 4].ToFloat(), + (clusterer.mPmemory->fragment).start + static_cast(peak.time()) + clustererNN.mOutputDataReg2_16[model_output_index + 2].ToFloat(), + clustererNN.mOutputDataReg2_16[model_output_index + 6].ToFloat(), + clustererNN.mClusterFlags[2 * glo_idx], + clustererNN.mClusterFlags[2 * glo_idx + 1]); + } else if (dtype == 1) { + pc.setFull(central_charge * clustererNN.mOutputDataReg2_32[model_output_index + 8], + static_cast(peak.pad()) + clustererNN.mOutputDataReg2_32[model_output_index], + clustererNN.mOutputDataReg2_32[model_output_index + 4], + (clusterer.mPmemory->fragment).start + static_cast(peak.time()) + clustererNN.mOutputDataReg2_32[model_output_index + 2], + clustererNN.mOutputDataReg2_32[model_output_index + 6], + clustererNN.mClusterFlags[2 * glo_idx], + clustererNN.mClusterFlags[2 * glo_idx + 1]); + } - uint32_t rowIndex = 0; - if (clusterOut != nullptr) { - rowIndex = GPUTPCCFClusterizer::sortIntoBuckets( - clusterer, - myCluster, - peak.row(), - clusterer.mNMaxClusterPerRow, - clusterer.mPclusterInRow, - clusterOut); - if (clusterer.mPclusterPosInRow != nullptr) { - clusterer.mPclusterPosInRow[full_glo_idx] = rowIndex; - } - } else if (clusterer.mPclusterPosInRow) { - rowIndex = clusterer.mPclusterPosInRow[full_glo_idx]; + tpc::ClusterNative myCluster; + bool rejectCluster = !pc.toNative(peak, central_charge, myCluster, clusterer.Param(), chargeMap); + if (clustererNN.mNnClusterizerUseClassification) { + rejectCluster |= (clustererNN.mOutputDataClass[CAMath::Min(full_glo_idx, (uint32_t)clusterer.mPmemory->counters.nClusters - 1)] <= 0); + } + if (rejectCluster) { + if (clusterer.mPclusterPosInRow) { + clusterer.mPclusterPosInRow[full_glo_idx] = clusterer.mNMaxClusterPerRow; } - CPU_ONLY(labelAcc->commit(peak.row(), rowIndex, clusterer.mNMaxClusterPerRow)); + return; + } - // Cluster 2 - if (dtype == 0) { - pc.setFull(central_charge * clustererNN.mOutputDataReg2_16[model_output_index + 9].ToFloat(), - static_cast(peak.pad()) + clustererNN.mOutputDataReg2_16[model_output_index + 1].ToFloat(), - clustererNN.mOutputDataReg2_16[model_output_index + 5].ToFloat(), - (clusterer.mPmemory->fragment).start + static_cast(peak.time()) + clustererNN.mOutputDataReg2_16[model_output_index + 3].ToFloat(), - clustererNN.mOutputDataReg2_16[model_output_index + 7].ToFloat(), - clustererNN.mClusterFlags[2 * glo_idx], - clustererNN.mClusterFlags[2 * glo_idx + 1]); - } else if (dtype == 1) { - pc.setFull(central_charge * clustererNN.mOutputDataReg2_32[model_output_index + 9], - static_cast(peak.pad()) + clustererNN.mOutputDataReg2_32[model_output_index + 1], - clustererNN.mOutputDataReg2_32[model_output_index + 5], - (clusterer.mPmemory->fragment).start + static_cast(peak.time()) + clustererNN.mOutputDataReg2_32[model_output_index + 3], - clustererNN.mOutputDataReg2_32[model_output_index + 7], - clustererNN.mClusterFlags[2 * glo_idx], - clustererNN.mClusterFlags[2 * glo_idx + 1]); + uint32_t rowIndex = 0; + if (clusterOut != nullptr) { + rowIndex = GPUTPCCFClusterizer::sortIntoBuckets( + clusterer, + myCluster, + peak.row(), + clusterer.mNMaxClusterPerRow, + clusterer.mPclusterInRow, + clusterOut); + if (clusterer.mPclusterPosInRow != nullptr) { + clusterer.mPclusterPosInRow[full_glo_idx] = rowIndex; } + } else if (clusterer.mPclusterPosInRow) { + rowIndex = clusterer.mPclusterPosInRow[full_glo_idx]; + } + CPU_ONLY(labelAcc->commit(peak.row(), rowIndex, clusterer.mNMaxClusterPerRow)); - rejectCluster = !pc.toNative(peak, central_charge, myCluster, clusterer.Param(), chargeMap); - if (rejectCluster) { - if (clusterer.mPclusterPosInRow) { - clusterer.mPclusterPosInRow[full_glo_idx] = clusterer.mNMaxClusterPerRow; - } - return; - } + // Cluster 2 + if (dtype == 0) { + pc.setFull(central_charge * clustererNN.mOutputDataReg2_16[model_output_index + 9].ToFloat(), + static_cast(peak.pad()) + clustererNN.mOutputDataReg2_16[model_output_index + 1].ToFloat(), + clustererNN.mOutputDataReg2_16[model_output_index + 5].ToFloat(), + (clusterer.mPmemory->fragment).start + static_cast(peak.time()) + clustererNN.mOutputDataReg2_16[model_output_index + 3].ToFloat(), + clustererNN.mOutputDataReg2_16[model_output_index + 7].ToFloat(), + clustererNN.mClusterFlags[2 * glo_idx], + clustererNN.mClusterFlags[2 * glo_idx + 1]); + } else if (dtype == 1) { + pc.setFull(central_charge * clustererNN.mOutputDataReg2_32[model_output_index + 9], + static_cast(peak.pad()) + clustererNN.mOutputDataReg2_32[model_output_index + 1], + clustererNN.mOutputDataReg2_32[model_output_index + 5], + (clusterer.mPmemory->fragment).start + static_cast(peak.time()) + clustererNN.mOutputDataReg2_32[model_output_index + 3], + clustererNN.mOutputDataReg2_32[model_output_index + 7], + clustererNN.mClusterFlags[2 * glo_idx], + clustererNN.mClusterFlags[2 * glo_idx + 1]); + } - if (clusterOut != nullptr) { - rowIndex = GPUTPCCFClusterizer::sortIntoBuckets( - clusterer, - myCluster, - peak.row(), - clusterer.mNMaxClusterPerRow, - clusterer.mPclusterInRow, - clusterOut); - if (clusterer.mPclusterPosInRow != nullptr) { - clusterer.mPclusterPosInRow[full_glo_idx] = rowIndex; - } - } else if (clusterer.mPclusterPosInRow) { - rowIndex = clusterer.mPclusterPosInRow[full_glo_idx]; - } - // CPU_ONLY(labelAcc->commit(peak.row(), rowIndex, clusterer.mNMaxClusterPerRow)); // -> Is this needed? How to handle MC labels for split clusters? - } else { + rejectCluster = !pc.toNative(peak, central_charge, myCluster, clusterer.Param(), chargeMap); + if (clustererNN.mNnClusterizerUseClassification) { + rejectCluster |= (clustererNN.mOutputDataClass[CAMath::Min(full_glo_idx, (uint32_t)clusterer.mPmemory->counters.nClusters - 1)] <= 0); + } + if (rejectCluster) { if (clusterer.mPclusterPosInRow) { clusterer.mPclusterPosInRow[full_glo_idx] = clusterer.mNMaxClusterPerRow; } return; } + + if (clusterOut != nullptr) { + rowIndex = GPUTPCCFClusterizer::sortIntoBuckets( + clusterer, + myCluster, + peak.row(), + clusterer.mNMaxClusterPerRow, + clusterer.mPclusterInRow, + clusterOut); + if (clusterer.mPclusterPosInRow != nullptr) { + clusterer.mPclusterPosInRow[full_glo_idx] = rowIndex; + } + } else if (clusterer.mPclusterPosInRow) { + rowIndex = clusterer.mPclusterPosInRow[full_glo_idx]; + } + // CPU_ONLY(labelAcc->commit(peak.row(), rowIndex, clusterer.mNMaxClusterPerRow)); // -> Is this needed? How to handle MC labels for split clusters? } // --------------------------------- @@ -579,27 +608,30 @@ template <> GPUdii() void GPUTPCNNClusterizerKernels::Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUSharedMemory& smem, processorType& processors, uint8_t sector, int8_t dtype, int8_t withMC, uint batchStart) { // Implements identical publishing logic as the heuristic clusterizer and deconvolution kernel - uint32_t idx = get_global_id(0); + uint32_t glo_idx = get_global_id(0); auto& clusterer = processors.tpcClusterer[sector]; auto& clustererNN = processors.tpcNNClusterer[sector]; + if (glo_idx + batchStart >= clusterer.mPmemory->counters.nClusters || glo_idx >= (uint32_t)clustererNN.mNnClusterizerBatchedMode) { + return; + } CfArray2D chargeMap(reinterpret_cast(clusterer.mPchargeMap)); - CfChargePos peak = clusterer.mPfilteredPeakPositions[idx + batchStart]; + CfChargePos peak = clusterer.mPfilteredPeakPositions[glo_idx + batchStart]; - clustererNN.mClusterFlags[2 * idx] = 0; - clustererNN.mClusterFlags[2 * idx + 1] = 0; + clustererNN.mClusterFlags[2 * glo_idx] = 0; + clustererNN.mClusterFlags[2 * glo_idx + 1] = 0; for (int i = 0; i < 8; i++) { Delta2 d = cfconsts::InnerNeighbors[i]; CfChargePos tmp_pos = peak.delta(d); PackedCharge charge = chargeMap[tmp_pos]; - clustererNN.mClusterFlags[2 * idx] += (d.y != 0 && charge.isSplit()); - clustererNN.mClusterFlags[2 * idx + 1] += (d.x != 0 && charge.isSplit()); + clustererNN.mClusterFlags[2 * glo_idx] += (d.y != 0 && charge.isSplit()); + clustererNN.mClusterFlags[2 * glo_idx + 1] += (d.x != 0 && charge.isSplit()); } for (int i = 0; i < 16; i++) { Delta2 d = cfconsts::OuterNeighbors[i]; CfChargePos tmp_pos = peak.delta(d); PackedCharge charge = chargeMap[tmp_pos]; - clustererNN.mClusterFlags[2 * idx] += (d.y != 0 && charge.isSplit() && !charge.has3x3Peak()); - clustererNN.mClusterFlags[2 * idx + 1] += (d.x != 0 && charge.isSplit() && !charge.has3x3Peak()); + clustererNN.mClusterFlags[2 * glo_idx] += (d.y != 0 && charge.isSplit() && !charge.has3x3Peak()); + clustererNN.mClusterFlags[2 * glo_idx + 1] += (d.x != 0 && charge.isSplit() && !charge.has3x3Peak()); } } @@ -607,7 +639,7 @@ GPUdii() void GPUTPCNNClusterizerKernels::Thread= o2::tpc::constants::MAXGLOBALPADROW) { - return 0; // Short-circuit for negative rows + return 0; // Short-circuit for out-of-bound rows } else { return (int)((GPUTPCGeometry::NPads(row_current) - GPUTPCGeometry::NPads(row_ref)) / 2); } @@ -623,11 +655,11 @@ GPUd() bool GPUTPCNNClusterizerKernels::isBoundary(int32_t row, int32_t pad, int if (pad < 0 || row < 0) { // Faster short-circuit return true; } else if (row < 63) { - return ((pad < 0) || (pad >= static_cast(GPUTPCGeometry::NPads(row)))); + return (pad >= static_cast(GPUTPCGeometry::NPads(row))); } else if (row < (63 + offset)) { // to account for the gap between IROC and OROC. Charge will be set to the boundary fill value in order to signal boundaries to the neural network return true; } else if (row < (o2::tpc::constants::MAXGLOBALPADROW + offset)) { - return ((pad < 0) || (pad >= static_cast(GPUTPCGeometry::NPads(row - offset)))); + return (pad >= static_cast(GPUTPCGeometry::NPads(row - offset))); } else { return true; } diff --git a/GPU/GPUTracking/TPCClusterFinder/GPUTPCNNClusterizerKernels.h b/GPU/GPUTracking/TPCClusterFinder/GPUTPCNNClusterizerKernels.h index 7469754594124..9c93726a097b7 100644 --- a/GPU/GPUTracking/TPCClusterFinder/GPUTPCNNClusterizerKernels.h +++ b/GPU/GPUTracking/TPCClusterFinder/GPUTPCNNClusterizerKernels.h @@ -15,12 +15,8 @@ #ifndef O2_GPU_NN_CLUSTERIZER_H #define O2_GPU_NN_CLUSTERIZER_H -#include "clusterFinderDefs.h" -#include "GPUGeneralKernels.h" -#include "GPUConstantMem.h" -#include "GPUTPCClusterFinder.h" #include "CfArray2D.h" -#include "PackedCharge.h" +#include "GPUGeneralKernels.h" #include "GPUTPCNNClusterizer.h" namespace o2::tpc @@ -33,6 +29,8 @@ namespace o2::gpu class ClusterAccumulator; class MCLabelAccumulator; +struct CfChargePos; +class PackedCharge; class GPUTPCNNClusterizerKernels : public GPUKernelTemplate { @@ -66,11 +64,6 @@ class GPUTPCNNClusterizerKernels : public GPUKernelTemplate GPUd() static void Thread(int32_t, int32_t, int32_t, int32_t, GPUSharedMemory&, processorType&, uint8_t = 0, int8_t = 0, int8_t = 0, uint = 0, Args...); private: - static GPUd() void fillInputData(int32_t, int32_t, int32_t, int32_t, processorType&, uint8_t, int8_t, uint); - static GPUd() void publishClustersReg1(uint, GPUSharedMemory&, processorType&, uint8_t, int8_t, int8_t, uint); - static GPUd() uint32_t sortIntoBuckets(GPUTPCClusterFinder&, const tpc::ClusterNative&, uint32_t, uint32_t, uint32_t*, tpc::ClusterNative*, uint32_t); - static GPUd() void publishClustersReg2(uint, GPUSharedMemory&, processorType&, uint8_t, int8_t, int8_t, uint); - static GPUd() int32_t padOffset(int32_t, int32_t); static GPUd() int32_t rowOffset(int32_t, int32_t); static GPUd() bool isBoundary(int32_t, int32_t, int32_t);