diff --git a/GPU/Common/GPUCommonAlgorithm.h b/GPU/Common/GPUCommonAlgorithm.h index ef9b54ef3f96d..428377527110a 100644 --- a/GPU/Common/GPUCommonAlgorithm.h +++ b/GPU/Common/GPUCommonAlgorithm.h @@ -41,7 +41,7 @@ class GPUCommonAlgorithm GPUd() static void sortInBlock(T* begin, T* end, const S& comp GPUCA_THREAD_INFO_DECL); template GPUd() static void sortDeviceDynamic(T* begin, T* end, const S& comp); -#if __cplusplus >= 202002L // sortOnDevice takes an auto parameter +#if __cplusplus >= 202002L // FIXME: auto parameter not supported in C++17 template GPUh() static void sortOnDevice(auto* rec, int32_t stream, T* begin, size_t N, const S& comp); #endif @@ -49,29 +49,26 @@ class GPUCommonAlgorithm GPUd() static void swap(T& a, T& b); private: - // Quicksort implementation template - GPUd() static void QuickSort(I f, I l) GPUnoexcept(); + GPUd() static void QuickSort(I f, I l) GPUnoexcept(); // Quicksort implementation - // Quicksort implementation template - GPUd() static void QuickSort(I f, I l, Cmp cmp) GPUnoexcept(); + GPUd() static void QuickSort(I f, I l, Cmp cmp) GPUnoexcept(); // Quicksort implementation + + template + GPUd() static void QuickSortImpl(I f, I l, Cmp cmp) GPUnoexcept(); // Quicksort implementation with index type IndexType, which must be able to hold the number of elements - // Insertionsort implementation template - GPUd() static void InsertionSort(I f, I l, Cmp cmp) GPUnoexcept(); + GPUd() static void InsertionSort(I f, I l, Cmp cmp) GPUnoexcept(); // Insertionsort implementation - // Helper for Quicksort implementation template - GPUd() static I MedianOf3Select(I f, I l, Cmp cmp) GPUnoexcept(); + GPUd() static I MedianOf3Select(I f, I l, Cmp cmp) GPUnoexcept(); // Helper for Quicksort implementation - // Helper for Quicksort implementation template - GPUd() static I UnguardedPartition(I f, I l, T piv, Cmp cmp) GPUnoexcept(); + GPUd() static I UnguardedPartition(I f, I l, T piv, Cmp cmp) GPUnoexcept(); // Helper for Quicksort implementation - // Helper template - GPUd() static void IterSwap(I a, I b) GPUnoexcept(); + GPUd() static void IterSwap(I a, I b) GPUnoexcept(); // Helper }; #ifndef GPUCA_ALGORITHM_STD @@ -147,11 +144,20 @@ GPUdi() I GPUCommonAlgorithm::UnguardedPartition(I f, I l, T piv, Cmp cmp) GPUno template GPUdi() void GPUCommonAlgorithm::QuickSort(I f, I l, Cmp cmp) GPUnoexcept() +{ + if (l - f <= 0xFFFF) { + QuickSortImpl(f, l, cmp); + } else { + QuickSortImpl(f, l, cmp); + } +} + +template +GPUdi() void GPUCommonAlgorithm::QuickSortImpl(I f, I l, Cmp cmp) GPUnoexcept() { if (f == l) { return; } - using IndexType = uint16_t; struct pair { IndexType first; @@ -159,7 +165,7 @@ GPUdi() void GPUCommonAlgorithm::QuickSort(I f, I l, Cmp cmp) GPUnoexcept() }; struct Stack { - pair data[11]; + pair data[sizeof(IndexType) * 8 - 4]; // The smaller partition is processed first, so the depth is bounded by log2(n / cutoff) + 1, cutoff = 50 > 2^5 uint8_t n{0}; GPUd() void emplace(IndexType x, IndexType y) @@ -282,7 +288,7 @@ GPUdi() void GPUCommonAlgorithm::sortInBlock(T* begin, T* end, const S& comp GPU { #ifndef GPUCA_GPUCODE GPUCommonAlgorithm::sort(begin, end, comp); -#elif defined(GPUCA_DETERMINISTIC_MODE) // Not using GPUCA_DETERMINISTIC_CODE, which is enforced in TPC compression +#elif defined(GPUCA_DETERMINISTIC_MODE) // Cannot use GPUCA_DETERMINISTIC_CODE macro here, which is enforced in TPC compression if (get_local_id(0) == 0) { GPUCommonAlgorithm::sort(begin, end, comp); } diff --git a/GPU/GPUTracking/Definitions/GPUSettingsList.h b/GPU/GPUTracking/Definitions/GPUSettingsList.h index 3c479646bf844..713168b8c3a5d 100644 --- a/GPU/GPUTracking/Definitions/GPUSettingsList.h +++ b/GPU/GPUTracking/Definitions/GPUSettingsList.h @@ -73,7 +73,7 @@ AddOptionRTC(tubeProtectMaxSize2, float, 2.f * 2.f, "", 0, "Square of max tube s AddOptionRTC(tubeProtectMinSize2, float, 0.5f * 0.5f, "", 0, "Square of min tube size (if larger than tubeProtectChi2)") AddOptionRTC(tubeRemoveSigma2, float, 1.25f * 1.25f, "", 0, "Max sigma2 to mark adjacent cluster for removal") AddOptionRTC(tubeRemoveMaxSize2, float, 2.5f * 2.5f, "", 0, "Square of max tube size (if smaller than tubeRejectChi2)") -AddOptionRTC(tubeExtraProtectMinOccupancy, uint32_t, 1500, "", 0, "Increase Protection, decrease removal by factor 2, when above this lokal occupancy / rowx") +AddOptionRTC(tubeExtraProtectMinOccupancy, uint32_t, 1500, "", 0, "Increase Protection, decrease removal by factor 2, when above this local occupancy / rowx") AddOptionRTC(clustersShiftTimebins, float, 0, "", 0, "Shift of TPC clusters (applied during CTF cluster decoding)") AddOptionRTC(clustersShiftTimebinsClusterizer, float, 0, "", 0, "Shift of TPC clusters (applied during CTF clusterization)") AddOptionRTC(clustersEdgeFixDistance, float, 0.f, "", 0, "If >0, revert cluster.flag edge bit distance to edge exceeds this parameter (fixed during CTF decoding)") @@ -113,7 +113,7 @@ AddOptionRTC(maxTimeBinAboveThresholdIn1000Bin, uint16_t, 500, "", 0, "Except pa 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") AddOptionRTC(hipTailFilter, uint8_t, 0, "", 0, "Enable Highly Ionising Particle tail filter in CheckPadBaseline (0 = disable, 1 = filter tails)") -AddOptionRTC(hipTailFilterMinimum, uint16_t, 1023, "", 0, "Thread signal above this minimum as saturated") +AddOptionRTC(hipTailFilterMinimum, uint16_t, 1023, "", 0, "Treat signal above this minimum as saturated") AddOptionRTC(hipTailFilterThreshold, uint16_t, 100, "", 0, "Threshold that must be exceeded for a timebin to be counted towards Highly Ionising Particle tail") AddOptionRTC(hipTailFilterAlpha, float, 0.5f, "", 0, "Smoothing factor for the exponential Highly Ionising Particle tail filter") AddOptionRTC(occupancyMapTimeBins, uint16_t, 16, "", 0, "Number of timebins per histogram bin of occupancy map (0 = disable occupancy map)") @@ -131,7 +131,7 @@ AddOptionRTC(trackFitMaxRowMissedHard, uint8_t, 10, "", 0, "Hard limit for numbe AddOptionRTC(extrapolationTrackingRowRange, uint8_t, 45, "", 0, "Number of rows from the upped/lower limit to search for extrapolated track candidates in for") AddOptionRTC(extrapolationTrackingMinRows, uint8_t, 10, "", 0, "Min num of rows an additional extrapolated track must span over") AddOptionRTC(extrapolationTrackingMinHits, uint8_t, 8, "", 0, "Min num of hits for an additional extrapolated track") -AddOptionRTC(noisyPadsQuickCheck, uint8_t, 1, "", 0, "Only check first fragment for noisy pads instead of all fragments (when test is enabled).") +AddOptionRTC(noisyPadsQuickCheck, uint8_t, 1, "", 0, "Only check first fragment for noisy pads instead of all fragments (when test is enabled)") AddOptionRTC(cfQMaxCutoff, uint8_t, 3, "", 0, "Cluster Finder rejects cluster with qmax below or equal to this threshold") AddOptionRTC(cfQTotCutoff, uint8_t, 5, "", 0, "Cluster Finder rejects cluster with qtot below or equal to this threshold") AddOptionRTC(cfQMaxCutoffSingleTime, uint8_t, 0, "", 0, "Cluster Finder rejects cluster with qMax below or equal to this threshold for single pad or single time bin clusters") @@ -209,7 +209,7 @@ BeginSubConfig(GPUSettingsRec, rec, configStandalone, "REC", 0, "Reconstruction AddOptionRTC(maxTrackQPtB5, float, 1.f / 0.010f, "", 0, "required max Q/Pt (==min Pt) of tracks") AddOptionRTC(fwdTPCDigitsAsClusters, uint8_t, 0, "", 0, "Forward TPC digits as clusters (if they pass the ZS threshold)") AddOptionRTC(bz0Pt10MeV, uint8_t, 60, "", 0, "Nominal Pt to set when bz = 0 (in 10 MeV)") -AddOptionRTC(fitInProjections, int8_t, -1, "", 0, "Fit in projection, -1 to enable full fit for all but passes but the first one") +AddOptionRTC(fitInProjections, int8_t, -1, "", 0, "Fit in projection, -1 fits all iterations but the last one in projections") AddOptionRTC(fitPropagateBzOnly, int8_t, -1, "", 0, "Propagate using Bz only for n passes") AddOptionRTC(useMatLUT, int8_t, 0, "", 0, "Use material lookup table for TPC refit") AddOptionRTC(trackingRefitGPUModel, int8_t, 1, "", 0, "Use GPU track model for the Global Track Refit") @@ -259,23 +259,23 @@ EndConfig() // Settings steering the processing of NN Clusterization BeginSubConfig(GPUSettingsProcessingNNclusterizer, nn, configStandalone.proc, "NN", 0, "Processing settings for neural network clusterizer", proc_nn) -AddOption(applyNNclusterizer, int, 0, "", 0, "(bool, default = 0), if the neural network clusterizer should be used.") +AddOption(applyNNclusterizer, int, 0, "", 0, "(bool, default = 0), if the neural network clusterizer should be used") AddOption(nnInferenceDevice, std::string, "CPU", "", 0, "(std::string) Specify inference device (cpu (default), rocm, cuda)") AddOption(nnInferenceDeviceId, unsigned int, 0, "", 0, "(unsigned int) Specify inference device id") AddOption(nnInferenceAllocateDevMem, int, 0, "", 0, "(bool, default = 0), if the device memory should be allocated for inference") AddOption(nnInferenceInputDType, std::string, "FP32", "", 0, "(std::string) Specify the datatype for which inference is performed (FP32: default, fp16)") // fp32 or fp16 AddOption(nnInferenceOutputDType, std::string, "FP32", "", 0, "(std::string) Specify the datatype for which inference is performed (fp32: default, fp16)") // fp32 or fp16 -AddOption(nnInferenceIntraOpNumThreads, int, 1, "", 0, "Number of threads used to evaluate one neural network (ONNX: SetIntraOpNumThreads). 0 = auto-detect, can lead to problems on SLURM systems.") -AddOption(nnInferenceInterOpNumThreads, int, 1, "", 0, "Number of threads used to evaluate one neural network (ONNX: SetInterOpNumThreads). 0 = auto-detect, can lead to problems on SLURM systems.") +AddOption(nnInferenceIntraOpNumThreads, int, 1, "", 0, "Number of threads used to evaluate one neural network (ONNX: SetIntraOpNumThreads). 0 = auto-detect, can lead to problems on SLURM systems") +AddOption(nnInferenceInterOpNumThreads, int, 1, "", 0, "Number of threads used to evaluate one neural network (ONNX: SetInterOpNumThreads). 0 = auto-detect, can lead to problems on SLURM systems") AddOption(nnInferenceEnableOrtOptimization, unsigned int, 99, "", 0, "Enables graph optimizations in ONNX Runtime. Can be [0, 1, 2, 99] -> see https://github.com/microsoft/onnxruntime/blob/3f71d637a83dc3540753a8bb06740f67e926dc13/include/onnxruntime/core/session/onnxruntime_c_api.h#L347") 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, 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") -AddOption(nnClusterizerSizeInputTime, int, 3, "", 0, "Size of the input to the NN (currently calcualted as (length-1)/2") +AddOption(nnClusterizerSizeInputRow, int, 3, "", 0, "Size of the input to the NN (currently calculated as (length-1)/2") +AddOption(nnClusterizerSizeInputPad, int, 3, "", 0, "Size of the input to the NN (currently calculated as (length-1)/2") +AddOption(nnClusterizerSizeInputTime, int, 3, "", 0, "Size of the input to the NN (currently calculated as (length-1)/2") 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") @@ -285,9 +285,9 @@ AddOption(nnClusterizerApplyNoiseSuppression, int, 1, "", 0, "Applies the NoiseS 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") AddOption(nnClassificationPath, std::string, "network_class.onnx", "", 0, "The classification network path") AddOption(nnRegressionPath, std::string, "network_reg.onnx", "", 0, "The regression network path") -AddOption(nnClassThreshold, float, 0.5, "", 0, "The cutoff at which clusters will be accepted / rejected.") -AddOption(nnSigmoidTrafoClassThreshold, int, 1, "", 0, "If true (default), then the classification threshold is transformed by an inverse sigmoid function. This depends on how the network was trained (with a sigmoid as acitvation function in the last layer or not).") -AddOption(nnEvalMode, std::string, "c1:r1", "", 0, "Concatention of modes, e.g. c1:r1 (classification class 1, regression class 1)") +AddOption(nnClassThreshold, float, 0.5, "", 0, "The cutoff at which clusters will be accepted / rejected") +AddOption(nnSigmoidTrafoClassThreshold, int, 1, "", 0, "If true (default), then the classification threshold is transformed by an inverse sigmoid function. This depends on how the network was trained (with a sigmoid as acitvation function in the last layer or not)") +AddOption(nnEvalMode, std::string, "c1:r1", "", 0, "Concatenation of modes, e.g. c1:r1 (classification class 1, regression class 1)") AddOption(nnClusterizerUseClassification, int, 1, "", 0, "If 1, the classification output of the network is used to select clusters, else only the regression output is used and no clusters are rejected by classification") AddOption(nnClusterizerForceGpuInputFill, int, 0, "", 0, "Forces to use the fillInputNNGPU function") // CCDB @@ -299,8 +299,8 @@ AddOption(nnCCDBWithMomentum, std::string, "", "", 0, "Distinguishes between the AddOption(nnCCDBClassificationLayerType, std::string, "FC", "", 0, "Distinguishes between network with different layer types. Options: FC, CNN") AddOption(nnCCDBRegressionLayerType, std::string, "FC", "", 0, "Distinguishes between network with different layer types. Options: FC, CNN") AddOption(nnCCDBBeamType, std::string, "pp", "", 0, "Distinguishes between networks trained for different beam types. Options: pp, pPb, PbPb") -AddOption(nnCCDBInteractionRate, std::string, "500", "", 0, "Distinguishes between networks for different interaction rates [kHz].") -AddOption(nnCCDBExtraMetadata, std::string, "", "", 0, "Extra metadata to distinguish between networks, e.g. for different internal datatypes, etc.") +AddOption(nnCCDBInteractionRate, std::string, "500", "", 0, "Distinguishes between networks for different interaction rates [kHz]") +AddOption(nnCCDBExtraMetadata, std::string, "", "", 0, "Extra metadata to distinguish between networks, e.g. for different internal datatypes, etc") AddHelp("help", 'h') EndConfig() @@ -341,8 +341,8 @@ AddOption(debugMask, uint32_t, (1 << 18) - 1, "debugMask", 0, "Mask for debug ou AddOption(debugLogSuffix, std::string, "", "debugSuffix", 0, "Suffix for debug log files with --debug 6") AddOption(debugDumpChecksum, bool, false, "", 0, "When using --debug 6, log only the checksum instead of writing the full dump to a file") AddOption(debugFileHexFloat, int32_t, -1, "", 0, "Use hex format to print floats to debug dump file") -AddOption(debugCSV, std::string, "", "", 0, "CSV filename to append the benchmark results. Verbosity determined by parameter --debug.") -AddOption(debugMarkdown, bool, false, "", 0, "Print the results of standlaone benchmarks in markdown format") +AddOption(debugCSV, std::string, "", "", 0, "CSV filename to append the benchmark results. Verbosity determined by parameter --debug") +AddOption(debugMarkdown, bool, false, "", 0, "Print the results of standalone benchmarks in markdown format") AddOption(serializeGPU, int8_t, 0, "", 0, "Synchronize after each kernel call (bit 1) and DMA transfer (bit 2) and identify failures") AddOption(recoTaskTiming, bool, 0, "", 0, "Perform summary timing after whole reconstruction tasks") AddOption(deterministicGPUReconstruction, int32_t, -1, "", 0, "Make CPU and GPU debug output comparable (sort / skip concurrent parts), -1 = automatic if debugLevel >= 6 or deterministic compile flag set", def(1)) @@ -353,8 +353,8 @@ AddOption(deviceTimers, bool, true, "", 0, "Use device timers instead of host-ba AddOption(keepAllMemory, bool, false, "", 0, "Allocate all memory on both device and host, and do not reuse") AddOption(keepDisplayMemory, bool, false, "", 0, "Like keepAllMemory, but only for memory required for event display") AddOption(disableMemoryReuse, bool, false, "", 0, "Disable memory reusage (for debugging only)") -AddOption(memoryAllocationStrategy, int8_t, 0, "", 0, "Memory Allocation Stragegy (0 = auto, 1 = individual allocations, 2 = single global allocation)") -AddOption(forceMemoryPoolSize, uint64_t, 1, "memSize", 0, "Force size of allocated GPU / page locked host memory", min(0ul)) +AddOption(memoryAllocationStrategy, int8_t, 0, "", 0, "Memory Allocation Strategy (0 = auto, 1 = individual allocations, 2 = single global allocation)") +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)) @@ -363,16 +363,16 @@ AddOption(forceMaxMemScalers, uint64_t, 0, "", 0, "Force using the maximum value AddOption(registerStandaloneInputMemory, bool, false, "registerInputMemory", 0, "Automatically register input memory buffers for the GPU") AddOption(nHostThreads, int32_t, -1, "nThreads", 't', "Number of host threads to run (-1: all)", min(-1), message("Using %s CPU threads")) AddOption(ompThreads, int32_t, -1, "", 0, "Deprecated synonym for nHostThreads") -AddOption(inKernelParallel, uint8_t, 2, "", 0, "Parallelize with multi-threading inside kernels on the host instead of over TPC sectors, 2 for nested parallelization over TPC sectors and inside kernels") +AddOption(inKernelParallel, uint8_t, 2, "", 0, "Parallelize with multi-threading inside kernels on the host instead of over TPC sectors: 0 disable, 1 full and exclusive in-kernel parallelization, 2 for nested parallelization over TPC sectors (nTPCClustererLanes thread-groups) and inside kernels") AddOption(autoAdjustHostThreads, bool, true, "", 0, "Auto-adjust number of OMP threads, decreasing the number for small input data") AddOption(nStreams, int8_t, 8, "", 0, "Number of GPU streams / command queues") AddOption(nTPCClustererLanes, int8_t, -1, "", 0, "Number of TPC clusterers that can run in parallel (-1 = autoset)") AddOption(overrideClusterizerFragmentLen, int32_t, -1, "", 0, "Force the cluster max fragment len to a certain value (-1 = autodetect)") AddOption(delayedOutput, bool, true, "", 0, "Delay output to be parallel to track fit") -AddOption(mergerSortTracks, int8_t, -1, "", 0, "Sort track indizes for GPU track fit") -AddOption(alternateBorderSort, int8_t, -1, "", 0, "Alternative implementation for sorting of border tracks") -AddOption(tpcCompressionGatherMode, int8_t, -1, "", 0, "TPC Compressed Clusters Gather Mode (0: DMA transfer gather gpu to host, 1: serial DMA to host and gather by copy on CPU, 2. gather via GPU kernal DMA access, 3. gather on GPU via kernel, dma afterwards") -AddOption(tpcCompressionGatherModeKernel, int8_t, -1, "", 0, "TPC Compressed Clusters Gather Mode Kernel (0: unbufferd, 1-3: buffered, 4: multi-block)") +AddOption(mergerSortTracks, int8_t, -1, "", 0, "Sort track indizes for GPU track fit (-1 = auto-device based on device parameters)") +AddOption(alternateBorderSort, int8_t, -1, "", 0, "Alternative implementation for sorting of border tracks (-1 = auto-device based on device parameters)") +AddOption(tpcCompressionGatherMode, int8_t, -1, "", 0, "TPC Compressed Clusters Gather Mode (0: DMA transfer gather gpu to host, 1: serial DMA to host and gather by copy on CPU, 2. gather via GPU kernal DMA access, 3. gather on GPU via kernel, dma afterwards, -1 = auto-device based on device parameters)") +AddOption(tpcCompressionGatherModeKernel, int8_t, -1, "", 0, "TPC Compressed Clusters Gather Mode Kernel (0: unbufferd, 1-3: buffered, 4: multi-block, -1 = auto-device based on device parameters)") AddOption(tpccfGatherKernel, bool, true, "", 0, "Use a kernel instead of the DMA engine to gather the clusters") AddOption(doublePipeline, bool, false, "", 0, "Double pipeline mode") AddOption(doublePipelineClusterizer, bool, true, "", 0, "Include the input data of the clusterizer in the double-pipeline") @@ -389,8 +389,8 @@ AddOption(tpcIncreasedMinClustersPerRow, uint32_t, 0, "", 0, "Impose a minimum b AddOption(noGPUMemoryRegistration, bool, false, "", 0, "Do not register input / output memory for GPU dma transfer") AddOption(o2PropagatorUseGPUField, bool, true, "", 0, "Makes the internal O2 propagator use the fast GPU polynomial b field approximation") AddOption(willProvideO2PropagatorLate, bool, false, "", 0, "Disable check for availability of o2 propagator and MatLUT at initialization") -AddOption(calibObjectsExtraMemorySize, uint32_t, 10u * 1024 * 1024, "", 0, "Extra spare memory added for calibration object buffer, to allow fow updates with larger objects") -AddOption(fastTransformObjectsMinMemorySize, uint32_t, 400u * 1024 * 1024, "", 0, "Extra spare memory added for calibration object buffer, to allow fow updates with larger objects") +AddOption(calibObjectsExtraMemorySize, uint32_t, 10u * 1024 * 1024, "", 0, "Extra spare memory added for calibration object buffer, to allow for updates with larger objects") +AddOption(fastTransformObjectsMinMemorySize, uint32_t, 400u * 1024 * 1024, "", 0, "Extra spare memory added for calibration object buffer, to allow for updates with larger objects") AddOption(lateO2MatLutProvisioningSize, uint32_t, 0u, "", 0, "Memory size to reserve for late provisioning of matlut table") AddOption(throttleAlarms, bool, false, "", 0, "Throttle rate at which alarms are sent to the InfoLogger in online runs") AddOption(outputSanityCheck, bool, false, "", 0, "Run some simple sanity checks finding errors in the output") @@ -543,7 +543,7 @@ AddOptionVec(compareInputNames, std::string, "QAinputName", 0, "Legend entries f AddOption(name, std::string, "", "", 0, "Legend entry for new data from current processing") AddOption(output, std::string, "", "", 0, "Store histograms in output root file", def(std::string("histograms.root"))) AddOption(inputHistogramsOnly, bool, false, "", 0, "Do not run tracking, but just create PDFs from input root files") -AddOption(strict, bool, true, "", 0, "Strict QA mode: Only consider resolution of tracks where the fit ended within 5 cm of the reference, and remove outliers.") +AddOption(strict, bool, true, "", 0, "Strict QA mode: Only consider resolution of tracks where the fit ended within 5 cm of the reference, and remove outliers") AddOption(qpt, float, 10.f, "", 0, "Set cut for Q/Pt", def(2.f)) AddOption(maxResX, float, 1e6f, "", 0, "Maxmimum X (~radius) for reconstructed track position to take into accound for resolution QA in cm") AddOption(recThreshold, float, 0.9f, "", 0, "Compute the efficiency including impure tracks with fake contamination") @@ -607,7 +607,7 @@ AddOption(runGPU, uint8_t, 1, "", 'g', "Use GPU for processing", message("GPU pr AddOptionSet(runGPU, uint8_t, 0, "", 'c', "Use CPU for processing", message("CPU enabled")) AddOption(gpuType, std::string, "AUTO", "", 0, "GPU type (CUDA / HIP / OCL / OCL) or CPU or AUTO") AddOption(runGPUforce, bool, true, "", 0, "Force usage of the specified GPU device type, no CPU fallback") -AddOption(noprompt, bool, true, "", 0, "Do prompt for keypress before exiting") +AddOption(promptOnExit, bool, false, "", 0, "Prompt for keypress before exiting") AddOption(continueOnError, bool, false, "", 0, "Continue processing after an error") AddOption(seed, int32_t, -1, "", 0, "Set srand seed (-1: random)") AddOption(StartEvent, int32_t, 0, "", 's', "First event to process", min(0)) @@ -643,11 +643,12 @@ AddOption(testSyncAsyncQcInSync, bool, false, "syncAsyncSyncQC", 0, "Run QC in s AddOption(timeFrameTime, bool, false, "tfTime", 0, "Print some debug information about time frame processing time") AddOption(controlProfiler, bool, false, "", 0, "Issues GPU profiler stop and start commands to profile only the relevant processing part") AddOption(preloadEvents, bool, false, "", 0, "Preload events into host memory before start processing") -AddOption(recoSteps, int32_t, -1, "", 0, "Bitmask for RecoSteps") -AddOption(recoStepsGPU, int32_t, -1, "", 0, "Bitmask for RecoSteps") +AddOption(recoSteps, int32_t, -1, "", 0, "Bitmask for RecoSteps. Run --listRecoSteps to get a list") +AddOption(recoStepsGPU, int32_t, -1, "", 0, "Bitmask for RecoSteps to run on GPU") +AddOption(listRecoSteps, bool, false, "", 0, "Print the list of RecoSteps with their bitmask values") AddOption(runMerger, int32_t, 1, "", 0, "Run track merging / refit", min(0), max(1)) -AddOption(runTRD, int32_t, -1, "", 0, "Enable TRD processing") -AddOption(rundEdx, int32_t, -1, "", 0, "Enable dEdx processing") +AddOption(runTRD, int32_t, -1, "", 0, "Enable TRD processing (-1 = decide automatically)") +AddOption(rundEdx, int32_t, -1, "", 0, "Enable dEdx processing (-1 = decide automatically)") AddOption(runCompression, int32_t, 1, "", 0, "Enable TPC Compression") AddOption(runTransformation, int32_t, 1, "", 0, "Enable TPC Transformation") AddOption(runRefit, bool, false, "", 0, "Enable final track refit") diff --git a/GPU/GPUTracking/Standalone/Benchmark/standalone.cxx b/GPU/GPUTracking/Standalone/Benchmark/standalone.cxx index 433668516dd58..5600191b1bc85 100644 --- a/GPU/GPUTracking/Standalone/Benchmark/standalone.cxx +++ b/GPU/GPUTracking/Standalone/Benchmark/standalone.cxx @@ -99,6 +99,12 @@ int32_t ReadConfiguration(int argc, char** argv) } return 1; } + if (configStandalone.listRecoSteps) { + printf("RecoSteps (bitmask for --recoSteps / --recoStepsGPU):\n"); + for (int32_t i = 0; i < gpudatatypes::N_RECO_STEPS; i++) { + printf("%10d (0x%04x): %s\n", 1 << i, 1 << i, gpudatatypes::RECO_STEP_NAMES[i]); + } + } if (configStandalone.printSettings > 1) { printf("Config Dump before ReadConfiguration\n"); qConfigPrint(); @@ -216,7 +222,7 @@ int32_t ReadConfiguration(int argc, char** argv) configStandalone.proc.outputSharedClusterMap = true; } if (configStandalone.eventDisplay) { - configStandalone.noprompt = 1; + configStandalone.promptOnExit = 0; } if (configStandalone.proc.debugLevel >= 4) { if (configStandalone.proc.inKernelParallel) { @@ -728,7 +734,7 @@ int32_t RunBenchmark(GPUReconstruction* recUse, GPUChainTracking* chainTrackingU if (tmpRetVal == GPUReconstruction::retValValue::retDoExit) { configStandalone.continueOnError = 0; // Forced exit from event display loop - configStandalone.noprompt = 1; + configStandalone.promptOnExit = 0; } if (tmpRetVal == GPUReconstruction::retValValue::retNonFatalErrorCode && configStandalone.proc.ignoreNonFatalGPUErrors) { printf("GPU Standalone Benchmark: Non-FATAL GPU error occured, ignoring\n"); @@ -1014,7 +1020,7 @@ int32_t main(int argc, char** argv) } rec->Exit(); - if (!configStandalone.noprompt) { + if (configStandalone.promptOnExit) { printf("Press a key to exit!\n"); getchar(); }