From db6effd041cc32abbb5d97b4a76d0b9a430b4431 Mon Sep 17 00:00:00 2001 From: Matteo Concas Date: Wed, 18 Sep 2024 09:52:43 +0200 Subject: [PATCH] Add processBulk to perform fits in parallel --- Common/DCAFitter/GPU/cuda/DCAFitterN.cu | 107 +++- .../GPU/cuda/test/testDCAFitterNGPU.cxx | 558 +++++++++++++++++- .../DCAFitter/include/DCAFitter/DCAFitterN.h | 16 +- Common/DCAFitter/test/testDCAFitterN.cxx | 54 +- 4 files changed, 660 insertions(+), 75 deletions(-) diff --git a/Common/DCAFitter/GPU/cuda/DCAFitterN.cu b/Common/DCAFitter/GPU/cuda/DCAFitterN.cu index 74c27796a6ebb..5cb93d474eb86 100644 --- a/Common/DCAFitter/GPU/cuda/DCAFitterN.cu +++ b/Common/DCAFitter/GPU/cuda/DCAFitterN.cu @@ -36,34 +36,51 @@ namespace o2::vertexing::device { namespace kernel { +GPUg() void warmUpGpuKernel() +{ + unsigned int tid = blockIdx.x * blockDim.x + threadIdx.x; + float ia, ib; + ia = ib = 0.0f; + ib += ia + tid; +} + template -GPUg() void printKernel(Fitter* ft) +GPUg() void printKernel(Fitter* fitter) { if (threadIdx.x == 0) { - printf(" =============== GPU DCA Fitter %d prongs ================\n", Fitter::getNProngs()); - ft->print(); + printf(" =============== GPU DCA Fitter %d prongs =================\n", Fitter::getNProngs()); + fitter->print(); printf(" =========================================================\n"); } } template -GPUg() void processKernel(Fitter* ft, int* res, Tr*... tracks) +GPUg() void processKernel(Fitter* fitter, int* res, Tr*... tracks) { - *res = ft->process(*tracks...); + *res = fitter->process(*tracks...); } + +template +GPUg() void processBulkKernel(Fitter* fitters, int* results, unsigned int N, Tr*... tracks) +{ + for (auto iThread{blockIdx.x * blockDim.x + threadIdx.x}; iThread < N; iThread += blockDim.x * gridDim.x) { + results[iThread] = fitters[iThread].process(tracks[iThread]...); + } +} + } // namespace kernel /// CPU handlers template void print(const int nBlocks, const int nThreads, - Fitter& ft) + Fitter& fitter) { - Fitter* ft_device; - gpuCheckError(cudaMalloc(reinterpret_cast(&ft_device), sizeof(Fitter))); - gpuCheckError(cudaMemcpy(ft_device, &ft, sizeof(Fitter), cudaMemcpyHostToDevice)); + Fitter* fitter_device; + gpuCheckError(cudaMalloc(reinterpret_cast(&fitter_device), sizeof(Fitter))); + gpuCheckError(cudaMemcpy(fitter_device, &fitter, sizeof(Fitter), cudaMemcpyHostToDevice)); - kernel::printKernel<<>>(ft_device); + kernel::printKernel<<>>(fitter_device); gpuCheckError(cudaPeekAtLastError()); gpuCheckError(cudaDeviceSynchronize()); @@ -75,11 +92,11 @@ int process(const int nBlocks, Fitter& fitter, Tr&... args) { - Fitter* ft_device; + Fitter* fitter_device; std::array tracks_device; int result, *result_device; - gpuCheckError(cudaMalloc(reinterpret_cast(&ft_device), sizeof(Fitter))); + gpuCheckError(cudaMalloc(reinterpret_cast(&fitter_device), sizeof(Fitter))); gpuCheckError(cudaMalloc(reinterpret_cast(&result_device), sizeof(int))); int iArg{0}; @@ -90,15 +107,15 @@ int process(const int nBlocks, }(), ...); - gpuCheckError(cudaMemcpy(ft_device, &fitter, sizeof(Fitter), cudaMemcpyHostToDevice)); + gpuCheckError(cudaMemcpy(fitter_device, &fitter, sizeof(Fitter), cudaMemcpyHostToDevice)); - std::apply([&](auto&&... args) { kernel::processKernel<<>>(ft_device, result_device, args...); }, tracks_device); + std::apply([&](auto&&... args) { kernel::processKernel<<>>(fitter_device, result_device, args...); }, tracks_device); gpuCheckError(cudaPeekAtLastError()); gpuCheckError(cudaDeviceSynchronize()); gpuCheckError(cudaMemcpy(&result, result_device, sizeof(int), cudaMemcpyDeviceToHost)); - gpuCheckError(cudaMemcpy(&fitter, ft_device, sizeof(Fitter), cudaMemcpyDeviceToHost)); + gpuCheckError(cudaMemcpy(&fitter, fitter_device, sizeof(Fitter), cudaMemcpyDeviceToHost)); iArg = 0; ([&] { gpuCheckError(cudaMemcpy(&args, tracks_device[iArg], sizeof(o2::track::TrackParCov), cudaMemcpyDeviceToHost)); @@ -107,11 +124,71 @@ int process(const int nBlocks, }(), ...); + gpuCheckError(cudaFree(fitter_device)); gpuCheckError(cudaFree(result_device)); return result; } +template +std::vector processBulk(const int nBlocks, + const int nThreads, + std::vector& fitters, + std::vector&... args) +{ + kernel::warmUpGpuKernel<<<1, 1>>>(); + + cudaEvent_t start, stop; + gpuCheckError(cudaEventCreate(&start)); + gpuCheckError(cudaEventCreate(&stop)); + const auto nFits{fitters.size()}; // for clarity: size of all the vectors needs to be equal, not enforcing it here yet. + std::vector results(nFits); + int* results_device; + Fitter* fitters_device; + std::array tracks_device; + + int iArg{0}; + ([&] { + gpuCheckError(cudaMalloc(reinterpret_cast(&(tracks_device[iArg])), sizeof(Tr) * args.size())); + gpuCheckError(cudaMemcpy(tracks_device[iArg], args.data(), sizeof(Tr) * args.size(), cudaMemcpyHostToDevice)); + ++iArg; + }(), + ...); + gpuCheckError(cudaMalloc(reinterpret_cast(&results_device), sizeof(int) * nFits)); + gpuCheckError(cudaMalloc(reinterpret_cast(&fitters_device), sizeof(Fitter) * nFits)); + gpuCheckError(cudaMemcpy(fitters_device, fitters.data(), sizeof(Fitter) * nFits, cudaMemcpyHostToDevice)); + + gpuCheckError(cudaEventRecord(start)); + std::apply([&](auto&&... args) { kernel::processBulkKernel<<>>(fitters_device, results_device, nFits, args...); }, tracks_device); + gpuCheckError(cudaEventRecord(stop)); + + gpuCheckError(cudaPeekAtLastError()); + gpuCheckError(cudaDeviceSynchronize()); + + gpuCheckError(cudaMemcpy(results.data(), results_device, sizeof(int) * results.size(), cudaMemcpyDeviceToHost)); + gpuCheckError(cudaMemcpy(fitters.data(), fitters_device, sizeof(Fitter) * nFits, cudaMemcpyDeviceToHost)); + + iArg = 0; + ([&] { + gpuCheckError(cudaMemcpy(args.data(), tracks_device[iArg], sizeof(Tr) * args.size(), cudaMemcpyDeviceToHost)); + gpuCheckError(cudaFree(tracks_device[iArg])); + ++iArg; + }(), + ...); + + gpuCheckError(cudaFree(fitters_device)); + gpuCheckError(cudaFree(results_device)); + gpuCheckError(cudaEventSynchronize(stop)); + + float milliseconds = 0; + gpuCheckError(cudaEventElapsedTime(&milliseconds, start, stop)); + + LOGP(info, "Kernel run in: {} ms using {} blocks and {} threads.", milliseconds, nBlocks, nThreads); + return results; +} + +template std::vector processBulk(const int, const int, std::vector>&, std::vector&, std::vector&); +template std::vector processBulk(const int, const int, std::vector>&, std::vector&, std::vector&, std::vector&); template int process(const int, const int, o2::vertexing::DCAFitterN<2>&, o2::track::TrackParCov&, o2::track::TrackParCov&); template int process(const int, const int, o2::vertexing::DCAFitterN<3>&, o2::track::TrackParCov&, o2::track::TrackParCov&, o2::track::TrackParCov&); template void print(const int, const int, o2::vertexing::DCAFitterN<2>&); diff --git a/Common/DCAFitter/GPU/cuda/test/testDCAFitterNGPU.cxx b/Common/DCAFitter/GPU/cuda/test/testDCAFitterNGPU.cxx index 8a2a70f65ac1c..ff4f12827f1e8 100644 --- a/Common/DCAFitter/GPU/cuda/test/testDCAFitterNGPU.cxx +++ b/Common/DCAFitter/GPU/cuda/test/testDCAFitterNGPU.cxx @@ -23,6 +23,10 @@ #include #include +#define nBlocks 60 +#define nThreads 1024 +#define NTest 1000000 + namespace o2 { namespace vertexing @@ -142,11 +146,11 @@ TLorentzVector generate(Vec3D& vtx, std::vector& vctr, f return parent; } +#ifdef DO_SINGLE_THREAD_TEST BOOST_AUTO_TEST_CASE(DCAFitterNProngs) { - gRandom->Delete(); - gRandom = new TRandom(42); - constexpr int NTest = 10000; + // gRandom->Delete(); + // gRandom = new TRandom(42); o2::utils::TreeStreamRedirector outStream("dcafitterNTest.root"); TGenPhaseSpace genPHS; @@ -222,17 +226,17 @@ BOOST_AUTO_TEST_CASE(DCAFitterNProngs) nfoundW++; } } - device::print(1, 1, ft); + meanDA /= nfoundA ? nfoundA : 1; - meanDAW /= nfoundA ? nfoundA : 1; + meanDAW /= nfoundAW ? nfoundA : 1; meanDW /= nfoundW ? nfoundW : 1; LOG(info) << "Processed " << NTest << " 2-prong vertices Helix : Helix"; LOG(info) << "2-prongs with abs.dist minization: eff= " << float(nfoundA) / NTest - << " mean.dist to truth: " << meanDA << " GPU time: " << swA.CpuTime(); + << " mean.dist to truth: " << meanDA << " Total time: " << swA.CpuTime() * 1000 << " ms"; LOG(info) << "2-prongs with abs.dist but wghPCA: eff= " << float(nfoundAW) / NTest - << " mean.dist to truth: " << meanDAW << " GPU time: " << swAW.CpuTime(); + << " mean.dist to truth: " << meanDAW << " Total time: " << swAW.CpuTime() * 1000 << " ms"; LOG(info) << "2-prongs with wgh.dist minization: eff= " << float(nfoundW) / NTest - << " mean.dist to truth: " << meanDW << " GPU time: " << swW.CpuTime(); + << " mean.dist to truth: " << meanDW << " Total time: " << swW.CpuTime() * 1000 << " ms"; BOOST_CHECK(nfoundA > 0.99 * NTest); BOOST_CHECK(nfoundAW > 0.99 * NTest); BOOST_CHECK(nfoundW > 0.99 * NTest); @@ -301,17 +305,16 @@ BOOST_AUTO_TEST_CASE(DCAFitterNProngs) } } - device::print(1, 1, ft); meanDA /= nfoundA ? nfoundA : 1; meanDAW /= nfoundA ? nfoundA : 1; meanDW /= nfoundW ? nfoundW : 1; LOG(info) << "Processed " << NTest << " 2-prong vertices Helix : Helix from gamma conversion"; LOG(info) << "2-prongs with abs.dist minization: eff= " << float(nfoundA) / NTest - << " mean.dist to truth: " << meanDA << " CPU time: " << swA.CpuTime(); + << " mean.dist to truth: " << meanDA << " Total time: " << swA.CpuTime(); LOG(info) << "2-prongs with abs.dist but wghPCA: eff= " << float(nfoundAW) / NTest - << " mean.dist to truth: " << meanDAW << " CPU time: " << swAW.CpuTime(); + << " mean.dist to truth: " << meanDAW << " Total time: " << swAW.CpuTime(); LOG(info) << "2-prongs with wgh.dist minization: eff= " << float(nfoundW) / NTest - << " mean.dist to truth: " << meanDW << " CPU time: " << swW.CpuTime(); + << " mean.dist to truth: " << meanDW << " Total time: " << swW.CpuTime(); BOOST_CHECK(nfoundA > 0.99 * NTest); BOOST_CHECK(nfoundAW > 0.99 * NTest); BOOST_CHECK(nfoundW > 0.99 * NTest); @@ -379,17 +382,17 @@ BOOST_AUTO_TEST_CASE(DCAFitterNProngs) nfoundW++; } } - device::print(1, 1, ft); + meanDA /= nfoundA ? nfoundA : 1; meanDAW /= nfoundAW ? nfoundAW : 1; meanDW /= nfoundW ? nfoundW : 1; LOG(info) << "Processed " << NTest << " 2-prong vertices: Helix : Line"; LOG(info) << "2-prongs with abs.dist minization: eff= " << float(nfoundA) / NTest - << " mean.dist to truth: " << meanDA << " CPU time: " << swA.CpuTime(); + << " mean.dist to truth: " << meanDA << " Total time: " << swA.CpuTime(); LOG(info) << "2-prongs with abs.dist but wghPCA: eff= " << float(nfoundAW) / NTest - << " mean.dist to truth: " << meanDAW << " CPU time: " << swAW.CpuTime(); + << " mean.dist to truth: " << meanDAW << " Total time: " << swAW.CpuTime(); LOG(info) << "2-prongs with wgh.dist minization: eff= " << float(nfoundW) / NTest - << " mean.dist to truth: " << meanDW << " CPU time: " << swW.CpuTime(); + << " mean.dist to truth: " << meanDW << " Total time: " << swW.CpuTime(); BOOST_CHECK(nfoundA > 0.99 * NTest); BOOST_CHECK(nfoundAW > 0.99 * NTest); BOOST_CHECK(nfoundW > 0.99 * NTest); @@ -456,17 +459,17 @@ BOOST_AUTO_TEST_CASE(DCAFitterNProngs) nfoundW++; } } - device::print(1, 1, ft); + meanDA /= nfoundA ? nfoundA : 1; meanDAW /= nfoundAW ? nfoundAW : 1; meanDW /= nfoundW ? nfoundW : 1; LOG(info) << "Processed " << NTest << " 2-prong vertices: Line : Line"; LOG(info) << "2-prongs with abs.dist minization: eff= " << float(nfoundA) / NTest - << " mean.dist to truth: " << meanDA << " CPU time: " << swA.CpuTime(); + << " mean.dist to truth: " << meanDA << " Total time: " << swA.CpuTime(); LOG(info) << "2-prongs with abs.dist but wghPCA: eff= " << float(nfoundAW) / NTest - << " mean.dist to truth: " << meanDAW << " CPU time: " << swAW.CpuTime(); + << " mean.dist to truth: " << meanDAW << " Total time: " << swAW.CpuTime(); LOG(info) << "2-prongs with wgh.dist minization: eff= " << float(nfoundW) / NTest - << " mean.dist to truth: " << meanDW << " CPU time: " << swW.CpuTime(); + << " mean.dist to truth: " << meanDW << " Total time: " << swW.CpuTime(); BOOST_CHECK(nfoundA > 0.99 * NTest); BOOST_CHECK(nfoundAW > 0.99 * NTest); BOOST_CHECK(nfoundW > 0.99 * NTest); @@ -532,17 +535,17 @@ BOOST_AUTO_TEST_CASE(DCAFitterNProngs) nfoundW++; } } - device::print(1, 1, ft); + meanDA /= nfoundA ? nfoundA : 1; meanDAW /= nfoundAW ? nfoundAW : 1; meanDW /= nfoundW ? nfoundW : 1; LOG(info) << "Processed " << NTest << " 3-prong vertices"; LOG(info) << "3-prongs with abs.dist minization: eff= " << float(nfoundA) / NTest - << " mean.dist to truth: " << meanDA << " CPU time: " << swA.CpuTime(); + << " mean.dist to truth: " << meanDA << " Total time: " << swA.CpuTime(); LOG(info) << "3-prongs with abs.dist but wghPCA: eff= " << float(nfoundAW) / NTest - << " mean.dist to truth: " << meanDAW << " CPU time: " << swAW.CpuTime(); + << " mean.dist to truth: " << meanDAW << " Total time: " << swAW.CpuTime(); LOG(info) << "3-prongs with wgh.dist minization: eff= " << float(nfoundW) / NTest - << " mean.dist to truth: " << meanDW << " CPU time: " << swW.CpuTime(); + << " mean.dist to truth: " << meanDW << " Total time: " << swW.CpuTime(); BOOST_CHECK(nfoundA > 0.99 * NTest); BOOST_CHECK(nfoundAW > 0.99 * NTest); BOOST_CHECK(nfoundW > 0.99 * NTest); @@ -553,6 +556,513 @@ BOOST_AUTO_TEST_CASE(DCAFitterNProngs) outStream.Close(); } +#endif + +BOOST_AUTO_TEST_CASE(DCAFitterNProngsBulk) +{ + // gRandom->Delete(); + // gRandom = new TRandom(42); + o2::utils::TreeStreamRedirector outStreamB("dcafitterNTestBulk.root"); + + TGenPhaseSpace genPHS; + constexpr double ele = 0.00051; + constexpr double gamma = 2 * ele + 1e-6; + constexpr double pion = 0.13957; + constexpr double k0 = 0.49761; + constexpr double kch = 0.49368; + constexpr double dch = 1.86965; + std::vector gammadec = {ele, ele}; + std::vector k0dec = {pion, pion}; + std::vector dchdec = {pion, kch, pion}; + std::vector> vctracks(3, std::vector(NTest)); + std::vector vtxGen(NTest); + + double bz = 5.0; + { // 2 prongs vertices bulk processing + LOG(info) << "\n\nBulk-processing 2-prong Helix - Helix case"; + std::vector forceQ{1, 1}; + + o2::vertexing::DCAFitterN<2> ft; // 2 prong fitter + ft.setBz(bz); + ft.setPropagateToPCA(true); // After finding the vertex, propagate tracks to the DCA. This is default anyway + ft.setMaxR(200); // do not consider V0 seeds with 2D circles crossing above this R. This is default anyway + ft.setMaxDZIni(4); // do not consider V0 seeds with tracks Z-distance exceeding this. This is default anyway + ft.setMaxDXYIni(4); // do not consider V0 seeds with tracks XY-distance exceeding this. This is default anyway + ft.setMinParamChange(1e-3); // stop iterations if max correction is below this value. This is default anyway + ft.setMinRelChi2Change(0.9); // stop iterations if chi2 improves by less that this factor + + std::vector> fitters_host(NTest); + std::vector genParents(NTest); + + std::string treeName2Abulk = "pr2aBulk", treeName2AWbulk = "pr2awBulk", treeName2Wbulk = "pr2wBulk"; + TStopwatch swAb, swAWb, swWb; + int nfoundAb = 0, nfoundAWb = 0, nfoundWb = 0; + double meanDAb = 0, meanDAWb = 0, meanDWb = 0; + swAb.Stop(); + swAWb.Stop(); + swWb.Stop(); + + ft.setUseAbsDCA(true); + std::fill(fitters_host.begin(), fitters_host.end(), ft); + for (int iev = 0; iev < NTest; iev++) { + std::vector vc(2); + genParents[iev] = generate(vtxGen[iev], vc, bz, genPHS, k0, k0dec, forceQ); + vctracks[0][iev] = vc[0]; + vctracks[1][iev] = vc[1]; + } + + swAb.Start(false); + auto ncAb = device::processBulk(nBlocks, nThreads, fitters_host, vctracks[0], vctracks[1]); // HERE WE FIT THE VERTICES + swAb.Stop(); + + for (int iev = 0; iev < NTest; iev++) { + LOG(debug) << "fit abs.dist " << iev << " NC: " << ncAb[iev] << " Chi2: " << (ncAb[iev] ? fitters_host[iev].getChi2AtPCACandidate(0) : -1); + if (ncAb[iev]) { + auto minDb = checkResults(outStreamB, treeName2Abulk, fitters_host[iev], vtxGen[iev], genParents[iev], k0dec); + meanDAb += minDb; + nfoundAb++; + } + } + + ft.setUseAbsDCA(true); + ft.setWeightedFinalPCA(true); + std::fill(fitters_host.begin(), fitters_host.end(), ft); + swAWb.Start(false); + auto ncAWb = device::processBulk(nBlocks, nThreads, fitters_host, vctracks[0], vctracks[1]); // HERE WE FIT THE VERTICES + swAWb.Stop(); + + for (int iev = 0; iev < NTest; iev++) { + LOG(debug) << "fit abs.dist with final weighted DCA " << iev << " NC: " << ncAWb[iev] << " Chi2: " << (ncAWb[iev] ? fitters_host[iev].getChi2AtPCACandidate(0) : -1); + if (ncAWb[iev]) { + auto minDb = checkResults(outStreamB, treeName2AWbulk, fitters_host[iev], vtxGen[iev], genParents[iev], k0dec); + meanDAWb += minDb; + nfoundAWb++; + } + } + + ft.setUseAbsDCA(false); + ft.setWeightedFinalPCA(false); + std::fill(fitters_host.begin(), fitters_host.end(), ft); + swWb.Start(false); + auto ncWb = device::processBulk(nBlocks, nThreads, fitters_host, vctracks[0], vctracks[1]); // HERE WE FIT THE VERTICES + swWb.Stop(); + + for (int iev = 0; iev < NTest; iev++) { + LOG(debug) << "fit wgh.dist " << iev << " NC: " << ncWb[iev] << " Chi2: " << (ncWb[iev] ? fitters_host[iev].getChi2AtPCACandidate(0) : -1); + if (ncWb[iev]) { + auto minDb = checkResults(outStreamB, treeName2Wbulk, fitters_host[iev], vtxGen[iev], genParents[iev], k0dec); + meanDWb += minDb; + nfoundWb++; + } + } + // + meanDAb /= nfoundAb ? nfoundAb : 1; + meanDAWb /= nfoundAWb ? nfoundAWb : 1; + meanDWb /= nfoundWb ? nfoundWb : 1; + LOGP(info, "Bulk-processed {} 2-prong vertices Helix : Helix", NTest); + LOG(info) << "2-prongs with abs.dist minization: eff= " << float(nfoundAb) / NTest + << " mean.dist to truth: " << meanDAb << " Total time: " << swAb.CpuTime() * 1000 << " ms"; + LOG(info) << "2-prongs with abs.dist but wghPCA: eff= " << float(nfoundAWb) / NTest + << " mean.dist to truth: " << meanDAWb << " Total time: " << swAWb.CpuTime() * 1000 << " ms"; + LOG(info) << "2-prongs with wgh.dist minization: eff= " << float(nfoundWb) / NTest + << " mean.dist to truth: " << meanDWb << " Total time: " << swWb.CpuTime() * 1000 << " ms"; + BOOST_CHECK(nfoundAb > 0.99 * NTest); + BOOST_CHECK(nfoundAWb > 0.99 * NTest); + BOOST_CHECK(nfoundWb > 0.99 * NTest); + BOOST_CHECK(meanDAb < 0.1); + BOOST_CHECK(meanDAWb < 0.1); + BOOST_CHECK(meanDWb < 0.1); + } + + { // 2 prongs vertices bulk processing for gamma conversion + LOG(info) << "\n\nBulk-processing 2-prong Helix - Helix case gamma conversion"; + std::vector forceQ{1, 1}; + + o2::vertexing::DCAFitterN<2> ft; // 2 prong fitter + ft.setBz(bz); + ft.setPropagateToPCA(true); // After finding the vertex, propagate tracks to the DCA. This is default anyway + ft.setMaxR(200); // do not consider V0 seeds with 2D circles crossing above this R. This is default anyway + ft.setMaxDZIni(4); // do not consider V0 seeds with tracks Z-distance exceeding this. This is default anyway + ft.setMaxDXYIni(4); // do not consider V0 seeds with tracks XY-distance exceeding this. This is default anyway + ft.setMinParamChange(1e-3); // stop iterations if max correction is below this value. This is default anyway + ft.setMinRelChi2Change(0.9); // stop iterations if chi2 improves by less that this factor + + std::vector> fitters_host(NTest); + std::vector genParents(NTest); + + std::string treeName2Abulk = "gpr2aBulk", treeName2AWbulk = "gpr2awBulk", treeName2Wbulk = "gpr2wBulk"; + TStopwatch swAb, swAWb, swWb; + int nfoundAb = 0, nfoundAWb = 0, nfoundWb = 0; + double meanDAb = 0, meanDAWb = 0, meanDWb = 0; + swAb.Stop(); + swAWb.Stop(); + swWb.Stop(); + + ft.setUseAbsDCA(true); + std::fill(fitters_host.begin(), fitters_host.end(), ft); + for (int iev = 0; iev < NTest; iev++) { + std::vector vc(2); + genParents[iev] = generate(vtxGen[iev], vc, bz, genPHS, gamma, gammadec, forceQ); + vctracks[0][iev] = vc[0]; + vctracks[1][iev] = vc[1]; + } + + swAb.Start(false); + auto ncAb = device::processBulk(nBlocks, nThreads, fitters_host, vctracks[0], vctracks[1]); // HERE WE FIT THE VERTICES + swAb.Stop(); + + for (int iev = 0; iev < NTest; iev++) { + LOG(debug) << "fit abs.dist " << iev << " NC: " << ncAb[iev] << " Chi2: " << (ncAb[iev] ? fitters_host[iev].getChi2AtPCACandidate(0) : -1); + if (ncAb[iev]) { + auto minDb = checkResults(outStreamB, treeName2Abulk, fitters_host[iev], vtxGen[iev], genParents[iev], gammadec); + meanDAb += minDb; + nfoundAb++; + } + } + // + ft.setUseAbsDCA(true); + ft.setWeightedFinalPCA(true); + std::fill(fitters_host.begin(), fitters_host.end(), ft); + swAWb.Start(false); + auto ncAWb = device::processBulk(nBlocks, nThreads, fitters_host, vctracks[0], vctracks[1]); // HERE WE FIT THE VERTICES + swAWb.Stop(); + + for (int iev = 0; iev < NTest; iev++) { + LOG(debug) << "fit abs.dist with final weighted DCA " << iev << " NC: " << ncAWb[iev] << " Chi2: " << (ncAWb[iev] ? fitters_host[iev].getChi2AtPCACandidate(0) : -1); + if (ncAWb[iev]) { + auto minDb = checkResults(outStreamB, treeName2AWbulk, fitters_host[iev], vtxGen[iev], genParents[iev], gammadec); + meanDAWb += minDb; + nfoundAWb++; + } + } + + ft.setUseAbsDCA(false); + ft.setWeightedFinalPCA(false); + std::fill(fitters_host.begin(), fitters_host.end(), ft); + swWb.Start(false); + auto ncWb = device::processBulk(nBlocks, nThreads, fitters_host, vctracks[0], vctracks[1]); // HERE WE FIT THE VERTICES + swWb.Stop(); + + for (int iev = 0; iev < NTest; iev++) { + LOG(debug) << "fit wgh.dist " << iev << " NC: " << ncWb[iev] << " Chi2: " << (ncWb[iev] ? fitters_host[iev].getChi2AtPCACandidate(0) : -1); + if (ncWb[iev]) { + auto minDb = checkResults(outStreamB, treeName2Wbulk, fitters_host[iev], vtxGen[iev], genParents[iev], gammadec); + meanDWb += minDb; + nfoundWb++; + } + } + // + + meanDAb /= nfoundAb ? nfoundAb : 1; + meanDAWb /= nfoundAWb ? nfoundAWb : 1; + meanDWb /= nfoundWb ? nfoundWb : 1; + LOGP(info, "Bulk-processed {} 2-prong vertices Helix : Helix from gamma conversion", NTest); + LOG(info) << "2-prongs with abs.dist minization: eff= " << float(nfoundAb) / NTest + << " mean.dist to truth: " << meanDAb << " Total time: " << swAb.CpuTime() * 1000 << " ms"; + LOG(info) << "2-prongs with abs.dist but wghPCA: eff= " << float(nfoundAWb) / NTest + << " mean.dist to truth: " << meanDAWb << " Total time: " << swAWb.CpuTime() * 1000 << " ms"; + LOG(info) << "2-prongs with wgh.dist minization: eff= " << float(nfoundWb) / NTest + << " mean.dist to truth: " << meanDWb << " Total time: " << swWb.CpuTime() * 1000 << " ms"; + BOOST_CHECK(nfoundAb > 0.99 * NTest); + BOOST_CHECK(nfoundAWb > 0.99 * NTest); + BOOST_CHECK(nfoundWb > 0.99 * NTest); + BOOST_CHECK(meanDAb < 2.1); + BOOST_CHECK(meanDAWb < 2.1); + BOOST_CHECK(meanDWb < 2.1); + } + + // 2 prongs vertices bulk processing with one of charges set to 0: Helix : Line + { + std::vector forceQ{1, 1}; + LOG(info) << "\n\nBulk-processing 2-prong Helix - Line case"; + o2::vertexing::DCAFitterN<2> ft; // 2 prong fitter + ft.setBz(bz); + ft.setPropagateToPCA(true); // After finding the vertex, propagate tracks to the DCA. This is default anyway + ft.setMaxR(200); // do not consider V0 seeds with 2D circles crossing above this R. This is default anyway + ft.setMaxDZIni(4); // do not consider V0 seeds with tracks Z-distance exceeding this. This is default anyway + ft.setMinParamChange(1e-3); // stop iterations if max correction is below this value. This is default anyway + ft.setMinRelChi2Change(0.9); // stop iterations if chi2 improves by less that this factor + + std::vector> fitters_host(NTest); + std::vector genParents(NTest); + + std::string treeName2Abulk = "pr2aHLb", treeName2AWbulk = "pr2awHLb", treeName2Wbulk = "pr2wHLb"; + TStopwatch swAb, swAWb, swWb; + int nfoundAb = 0, nfoundAWb = 0, nfoundWb = 0; + double meanDAb = 0, meanDAWb = 0, meanDWb = 0; + swAb.Stop(); + swAWb.Stop(); + swWb.Stop(); + + for (int iev = 0; iev < NTest; iev++) { + forceQ[iev % 2] = 1; + forceQ[1 - iev % 2] = 0; + std::vector vc(2); + genParents[iev] = generate(vtxGen[iev], vc, bz, genPHS, k0, k0dec, forceQ); + vctracks[0][iev] = vc[0]; + vctracks[1][iev] = vc[1]; + } + ft.setUseAbsDCA(true); + std::fill(fitters_host.begin(), fitters_host.end(), ft); + + swAb.Start(false); + auto ncAb = device::processBulk(nBlocks, nThreads, fitters_host, vctracks[0], vctracks[1]); // HERE WE FIT THE VERTICES + swAb.Stop(); + + for (int iev = 0; iev < NTest; iev++) { + LOG(debug) << "fit abs.dist with final weighted DCA " << iev << " NC: " << ncAb[iev] << " Chi2: " << (ncAb[iev] ? fitters_host[iev].getChi2AtPCACandidate(0) : -1); + if (ncAb[iev]) { + auto minDb = checkResults(outStreamB, treeName2Abulk, fitters_host[iev], vtxGen[iev], genParents[iev], k0dec); + meanDAb += minDb; + nfoundAb++; + } + } + + ft.setUseAbsDCA(true); + ft.setWeightedFinalPCA(true); + std::fill(fitters_host.begin(), fitters_host.end(), ft); + swAWb.Start(false); + auto ncAWb = device::processBulk(nBlocks, nThreads, fitters_host, vctracks[0], vctracks[1]); // HERE WE FIT THE VERTICES + swAWb.Stop(); + + for (int iev = 0; iev < NTest; iev++) { + LOG(debug) << "fit abs.dist " << iev << " NC: " << ncAWb[iev] << " Chi2: " << (ncAWb[iev] ? fitters_host[iev].getChi2AtPCACandidate(0) : -1); + if (ncAWb[iev]) { + auto minDb = checkResults(outStreamB, treeName2AWbulk, fitters_host[iev], vtxGen[iev], genParents[iev], k0dec); + meanDAWb += minDb; + nfoundAWb++; + } + } + + ft.setUseAbsDCA(false); + ft.setWeightedFinalPCA(false); + std::fill(fitters_host.begin(), fitters_host.end(), ft); + swWb.Start(false); + auto ncWb = device::processBulk(nBlocks, nThreads, fitters_host, vctracks[0], vctracks[1]); // HERE WE FIT THE VERTICES + swWb.Stop(); + + for (int iev = 0; iev < NTest; iev++) { + LOG(debug) << "fit wgh.dist " << iev << " NC: " << ncWb[iev] << " Chi2: " << (ncWb[iev] ? fitters_host[iev].getChi2AtPCACandidate(0) : -1); + if (ncWb[iev]) { + auto minDb = checkResults(outStreamB, treeName2Wbulk, fitters_host[iev], vtxGen[iev], genParents[iev], k0dec); + meanDWb += minDb; + nfoundWb++; + } + } + + // + meanDAb /= nfoundAb ? nfoundAb : 1; + meanDAWb /= nfoundAWb ? nfoundAWb : 1; + meanDWb /= nfoundWb ? nfoundWb : 1; + LOG(info) << "Bulk-processed " << NTest << " 2-prong vertices: Helix : Line"; + LOG(info) << "2-prongs with abs.dist minization: eff= " << float(nfoundAb) / NTest + << " mean.dist to truth: " << meanDAb << " Total time: " << swAb.CpuTime() * 1000 << " ms"; + LOG(info) << "2-prongs with abs.dist but wghPCA: eff= " << float(nfoundAWb) / NTest + << " mean.dist to truth: " << meanDAWb << " Total time: " << swAWb.CpuTime() * 1000 << " ms"; + LOG(info) << "2-prongs with wgh.dist minization: eff= " << float(nfoundWb) / NTest + << " mean.dist to truth: " << meanDWb << " Total time: " << swWb.CpuTime() * 1000 << " ms"; + BOOST_CHECK(nfoundAb > 0.99 * NTest); + BOOST_CHECK(nfoundAWb > 0.99 * NTest); + BOOST_CHECK(nfoundWb > 0.99 * NTest); + BOOST_CHECK(meanDAb < 0.1); + BOOST_CHECK(meanDAWb < 0.1); + BOOST_CHECK(meanDWb < 0.1); + } + + // 2 prongs vertices with both of charges set to 0: Line : Line + { + std::vector forceQ{0, 0}; + LOG(info) << "\n\nBulk-processing 2-prong Line - Line case"; + o2::vertexing::DCAFitterN<2> ft; // 2 prong fitter + ft.setBz(bz); + ft.setPropagateToPCA(true); // After finding the vertex, propagate tracks to the DCA. This is default anyway + ft.setMaxR(200); // do not consider V0 seeds with 2D circles crossing above this R. This is default anyway + ft.setMaxDZIni(4); // do not consider V0 seeds with tracks Z-distance exceeding this. This is default anyway + ft.setMinParamChange(1e-3); // stop iterations if max correction is below this value. This is default anyway + ft.setMinRelChi2Change(0.9); // stop iterations if chi2 improves by less that this factor + + std::vector> fitters_host(NTest); + std::vector genParents(NTest); + + std::string treeName2Abulk = "pr2aLL", treeName2AWbulk = "pr2awLL", treeName2Wbulk = "pr2wLL"; + TStopwatch swAb, swAWb, swWb; + int nfoundAb = 0, nfoundAWb = 0, nfoundWb = 0; + double meanDAb = 0, meanDAWb = 0, meanDWb = 0; + swAb.Stop(); + swAWb.Stop(); + swWb.Stop(); + for (int iev = 0; iev < NTest; iev++) { + forceQ[0] = forceQ[1] = 0; + std::vector vc(2); + genParents[iev] = generate(vtxGen[iev], vc, bz, genPHS, k0, k0dec, forceQ); + vctracks[0][iev] = vc[0]; + vctracks[1][iev] = vc[1]; + } + + ft.setUseAbsDCA(true); + std::fill(fitters_host.begin(), fitters_host.end(), ft); + + swAb.Start(false); + auto ncAb = device::processBulk(nBlocks, nThreads, fitters_host, vctracks[0], vctracks[1]); // HERE WE FIT THE VERTICES + swAb.Stop(); + + for (int iev = 0; iev < NTest; iev++) { + LOG(debug) << "fit abs.dist " << iev << " NC: " << ncAb[iev] << " Chi2: " << (ncAb[iev] ? fitters_host[iev].getChi2AtPCACandidate(0) : -1); + if (ncAb[iev]) { + auto minDb = checkResults(outStreamB, treeName2Abulk, fitters_host[iev], vtxGen[iev], genParents[iev], k0dec); + meanDAb += minDb; + nfoundAb++; + } + } + + ft.setUseAbsDCA(true); + ft.setWeightedFinalPCA(true); + std::fill(fitters_host.begin(), fitters_host.end(), ft); + swAWb.Start(false); + auto ncAWb = device::processBulk(nBlocks, nThreads, fitters_host, vctracks[0], vctracks[1]); // HERE WE FIT THE VERTICES + swAWb.Stop(); + for (int iev = 0; iev < NTest; iev++) { + LOG(debug) << "fit abs.dist " << iev << " NC: " << ncAWb[iev] << " Chi2: " << (ncAWb[iev] ? fitters_host[iev].getChi2AtPCACandidate(0) : -1); + if (ncAWb[iev]) { + auto minDb = checkResults(outStreamB, treeName2AWbulk, fitters_host[iev], vtxGen[iev], genParents[iev], k0dec); + meanDAWb += minDb; + nfoundAWb++; + } + } + + ft.setUseAbsDCA(false); + ft.setWeightedFinalPCA(false); + std::fill(fitters_host.begin(), fitters_host.end(), ft); + + swWb.Start(false); + auto ncWb = device::processBulk(nBlocks, nThreads, fitters_host, vctracks[0], vctracks[1]); // HERE WE FIT THE VERTICES + swWb.Stop(); + + for (int iev = 0; iev < NTest; iev++) { + LOG(debug) << "fit wgh.dist " << iev << " NC: " << ncWb[iev] << " Chi2: " << (ncWb[iev] ? fitters_host[iev].getChi2AtPCACandidate(0) : -1); + if (ncWb[iev]) { + auto minDb = checkResults(outStreamB, treeName2Wbulk, fitters_host[iev], vtxGen[iev], genParents[iev], k0dec); + meanDWb += minDb; + nfoundWb++; + } + } + // ft.print(); + meanDAb /= nfoundAb ? nfoundAb : 1; + meanDAWb /= nfoundAWb ? nfoundAWb : 1; + meanDWb /= nfoundWb ? nfoundWb : 1; + LOG(info) << "Bulk-processed " << NTest << " 2-prong vertices: Line : Line"; + LOG(info) << "2-prongs with abs.dist minization: eff= " << float(nfoundAb) / NTest + << " mean.dist to truth: " << meanDAb << " Total time: " << swAb.CpuTime() * 1000 << " ms"; + LOG(info) << "2-prongs with abs.dist but wghPCA: eff= " << float(nfoundAWb) / NTest + << " mean.dist to truth: " << meanDAWb << " Total time: " << swAWb.CpuTime() * 1000 << " ms"; + LOG(info) << "2-prongs with wgh.dist minization: eff= " << float(nfoundWb) / NTest + << " mean.dist to truth: " << meanDWb << " Total time: " << swWb.CpuTime() * 1000 << " ms"; + BOOST_CHECK(nfoundAb > 0.99 * NTest); + BOOST_CHECK(nfoundAWb > 0.99 * NTest); + BOOST_CHECK(nfoundWb > 0.99 * NTest); + BOOST_CHECK(meanDAb < 0.1); + BOOST_CHECK(meanDAWb < 0.1); + BOOST_CHECK(meanDWb < 0.1); + } + + // Bulk-process 3 prongs vertices + { + LOG(info) << "\n\nBulk-processing 3-prongs"; + std::vector forceQ{1, 1, 1}; + + o2::vertexing::DCAFitterN<3> ft; // 3 prong fitter + ft.setBz(bz); + ft.setPropagateToPCA(true); // After finding the vertex, propagate tracks to the DCA. This is default anyway + ft.setMaxR(200); // do not consider V0 seeds with 2D circles crossing above this R. This is default anyway + ft.setMaxDZIni(4); // do not consider V0 seeds with tracks Z-distance exceeding this. This is default anyway + ft.setMinParamChange(1e-3); // stop iterations if max correction is below this value. This is default anyway + ft.setMinRelChi2Change(0.9); // stop iterations if chi2 improves by less that this factor + + std::vector> fitters_host(NTest); + std::vector genParents(NTest); + + std::string treeName3Abulk = "pr3a", treeName3AWbulk = "pr3aw", treeName3Wbulk = "pr3w"; + TStopwatch swAb, swAWb, swWb; + int nfoundAb = 0, nfoundAWb = 0, nfoundWb = 0; + double meanDAb = 0, meanDAWb = 0, meanDWb = 0; + swAb.Stop(); + swAWb.Stop(); + swWb.Stop(); + for (int iev = 0; iev < NTest; iev++) { + std::vector vc(3); + genParents[iev] = generate(vtxGen[iev], vc, bz, genPHS, dch, dchdec, forceQ); + + vctracks[0][iev] = vc[0]; + vctracks[1][iev] = vc[1]; + vctracks[2][iev] = vc[2]; + } + + ft.setUseAbsDCA(true); + std::fill(fitters_host.begin(), fitters_host.end(), ft); + swAb.Start(false); + auto ncAb = device::processBulk(nBlocks, nThreads, fitters_host, vctracks[0], vctracks[1], vctracks[2]); // HERE WE FIT THE VERTICES + swAb.Stop(); + for (int iev = 0; iev < NTest; iev++) { + LOG(debug) << "fit abs.dist " << iev << " NC: " << ncAb[iev] << " Chi2: " << (ncAb[iev] ? fitters_host[iev].getChi2AtPCACandidate(0) : -1); + if (ncAb[iev]) { + auto minDb = checkResults(outStreamB, treeName3Abulk, fitters_host[iev], vtxGen[iev], genParents[iev], dchdec); + meanDAb += minDb; + nfoundAb++; + } + } + + ft.setUseAbsDCA(true); + ft.setWeightedFinalPCA(true); + std::fill(fitters_host.begin(), fitters_host.end(), ft); + + swAWb.Start(false); + auto ncAWb = device::processBulk(nBlocks, nThreads, fitters_host, vctracks[0], vctracks[1], vctracks[2]); // HERE WE FIT THE VERTICES + swAWb.Stop(); + for (int iev = 0; iev < NTest; iev++) { + LOG(debug) << "fit abs.dist " << iev << " NC: " << ncAWb[iev] << " Chi2: " << (ncAWb[iev] ? fitters_host[iev].getChi2AtPCACandidate(0) : -1); + if (ncAWb[iev]) { + auto minDb = checkResults(outStreamB, treeName3AWbulk, fitters_host[iev], vtxGen[iev], genParents[iev], dchdec); + meanDAWb += minDb; + nfoundAWb++; + } + } + + ft.setUseAbsDCA(false); + ft.setWeightedFinalPCA(false); + std::fill(fitters_host.begin(), fitters_host.end(), ft); + + swWb.Start(false); + auto ncWb = device::processBulk(nBlocks, nThreads, fitters_host, vctracks[0], vctracks[1], vctracks[2]); // HERE WE FIT THE VERTICES + swWb.Stop(); + for (int iev = 0; iev < NTest; iev++) { + LOG(debug) << "fit wgh.dist " << iev << " NC: " << ncWb[iev] << " Chi2: " << (ncWb[iev] ? fitters_host[iev].getChi2AtPCACandidate(0) : -1); + if (ncWb[iev]) { + auto minDb = checkResults(outStreamB, treeName3Wbulk, fitters_host[iev], vtxGen[iev], genParents[iev], dchdec); + meanDWb += minDb; + nfoundWb++; + } + } + + // ft.print(); + meanDAb /= nfoundAb ? nfoundAb : 1; + meanDAWb /= nfoundAWb ? nfoundAWb : 1; + meanDWb /= nfoundWb ? nfoundWb : 1; + LOG(info) << "Bulk-processed " << NTest << " 3-prong vertices"; + LOG(info) << "3-prongs with abs.dist minization: eff= " << float(nfoundAb) / NTest + << " mean.dist to truth: " << meanDAb << " Total time: " << swAb.CpuTime() * 1000 << " ms"; + LOG(info) << "3-prongs with abs.dist but wghPCA: eff= " << float(nfoundAWb) / NTest + << " mean.dist to truth: " << meanDAWb << " Total time: " << swAWb.CpuTime() * 1000 << " ms"; + LOG(info) << "3-prongs with wgh.dist minization: eff= " << float(nfoundWb) / NTest + << " mean.dist to truth: " << meanDWb << " Total time: " << swWb.CpuTime() * 1000 << " ms"; + BOOST_CHECK(nfoundAb > 0.99 * NTest); + BOOST_CHECK(nfoundAWb > 0.99 * NTest); + BOOST_CHECK(nfoundWb > 0.99 * NTest); + BOOST_CHECK(meanDAb < 0.1); + BOOST_CHECK(meanDAWb < 0.1); + BOOST_CHECK(meanDWb < 0.1); + } + outStreamB.Close(); +} } // namespace vertexing } // namespace o2 \ No newline at end of file diff --git a/Common/DCAFitter/include/DCAFitter/DCAFitterN.h b/Common/DCAFitter/include/DCAFitter/DCAFitterN.h index b063d1f587fe9..5a89597ad379a 100644 --- a/Common/DCAFitter/include/DCAFitter/DCAFitterN.h +++ b/Common/DCAFitter/include/DCAFitter/DCAFitterN.h @@ -17,10 +17,10 @@ #ifndef _ALICEO2_DCA_FITTERN_ #define _ALICEO2_DCA_FITTERN_ -#include "MathUtils/Cartesian.h" -#include "ReconstructionDataFormats/Track.h" #include "DCAFitter/HelixHelper.h" #include "DetectorsBase/Propagator.h" +#include "MathUtils/Cartesian.h" +#include "ReconstructionDataFormats/Track.h" namespace o2 { @@ -1136,15 +1136,13 @@ using DCAFitter3 = DCAFitterN<3, o2::track::TrackParCov>; namespace device { template -void print(const int nBlocks, - const int nThreads, - Fitter& ft); +void print(const int nBlocks, const int nThreads, Fitter& ft); template -int process(const int nBlocks, - const int nThreads, - Fitter&, - Tr&... args); +int process(const int nBlocks, const int nThreads, Fitter&, Tr&... args); + +template +std::vector processBulk(const int nBlocks, const int nThreads, std::vector& fitters, std::vector
&... args); } // namespace device } // namespace vertexing diff --git a/Common/DCAFitter/test/testDCAFitterN.cxx b/Common/DCAFitter/test/testDCAFitterN.cxx index bfd671315adf6..2f9c4d455376e 100644 --- a/Common/DCAFitter/test/testDCAFitterN.cxx +++ b/Common/DCAFitter/test/testDCAFitterN.cxx @@ -164,7 +164,7 @@ BOOST_AUTO_TEST_CASE(DCAFitterNProngs) double bz = 5.0; // 2 prongs vertices { - LOG(info) << "Processing 2-prong Helix - Helix case"; + LOG(info) << "\n\nProcessing 2-prong Helix - Helix case"; std::vector forceQ{1, 1}; o2::vertexing::DCAFitterN<2> ft; // 2 prong fitter @@ -221,17 +221,17 @@ BOOST_AUTO_TEST_CASE(DCAFitterNProngs) nfoundW++; } } - ft.print(); + // ft.print(); meanDA /= nfoundA ? nfoundA : 1; - meanDAW /= nfoundA ? nfoundA : 1; + meanDAW /= nfoundAW ? nfoundAW : 1; meanDW /= nfoundW ? nfoundW : 1; LOG(info) << "Processed " << NTest << " 2-prong vertices Helix : Helix"; LOG(info) << "2-prongs with abs.dist minization: eff= " << float(nfoundA) / NTest - << " mean.dist to truth: " << meanDA << " CPU time: " << swA.CpuTime(); + << " mean.dist to truth: " << meanDA << " CPU time: " << swA.CpuTime() * 1000 << " ms"; LOG(info) << "2-prongs with abs.dist but wghPCA: eff= " << float(nfoundAW) / NTest - << " mean.dist to truth: " << meanDAW << " CPU time: " << swAW.CpuTime(); + << " mean.dist to truth: " << meanDAW << " CPU time: " << swAW.CpuTime() * 1000 << " ms"; LOG(info) << "2-prongs with wgh.dist minization: eff= " << float(nfoundW) / NTest - << " mean.dist to truth: " << meanDW << " CPU time: " << swW.CpuTime(); + << " mean.dist to truth: " << meanDW << " CPU time: " << swW.CpuTime() * 1000 << " ms"; BOOST_CHECK(nfoundA > 0.99 * NTest); BOOST_CHECK(nfoundAW > 0.99 * NTest); BOOST_CHECK(nfoundW > 0.99 * NTest); @@ -242,7 +242,7 @@ BOOST_AUTO_TEST_CASE(DCAFitterNProngs) // 2 prongs vertices with collinear tracks (gamma conversion) { - LOG(info) << "Processing 2-prong Helix - Helix case gamma conversion"; + LOG(info) << "\n\nProcessing 2-prong Helix - Helix case gamma conversion"; std::vector forceQ{1, 1}; o2::vertexing::DCAFitterN<2> ft; // 2 prong fitter @@ -299,17 +299,17 @@ BOOST_AUTO_TEST_CASE(DCAFitterNProngs) nfoundW++; } } - ft.print(); + // ft.print(); meanDA /= nfoundA ? nfoundA : 1; - meanDAW /= nfoundA ? nfoundA : 1; + meanDAW /= nfoundAW ? nfoundAW : 1; meanDW /= nfoundW ? nfoundW : 1; LOG(info) << "Processed " << NTest << " 2-prong vertices Helix : Helix from gamma conversion"; LOG(info) << "2-prongs with abs.dist minization: eff= " << float(nfoundA) / NTest - << " mean.dist to truth: " << meanDA << " CPU time: " << swA.CpuTime(); + << " mean.dist to truth: " << meanDA << " CPU time: " << swA.CpuTime() * 1000 << " ms"; LOG(info) << "2-prongs with abs.dist but wghPCA: eff= " << float(nfoundAW) / NTest - << " mean.dist to truth: " << meanDAW << " CPU time: " << swAW.CpuTime(); + << " mean.dist to truth: " << meanDAW << " CPU time: " << swAW.CpuTime() * 1000 << " ms"; LOG(info) << "2-prongs with wgh.dist minization: eff= " << float(nfoundW) / NTest - << " mean.dist to truth: " << meanDW << " CPU time: " << swW.CpuTime(); + << " mean.dist to truth: " << meanDW << " CPU time: " << swW.CpuTime() * 1000 << " ms"; BOOST_CHECK(nfoundA > 0.99 * NTest); BOOST_CHECK(nfoundAW > 0.99 * NTest); BOOST_CHECK(nfoundW > 0.99 * NTest); @@ -321,7 +321,7 @@ BOOST_AUTO_TEST_CASE(DCAFitterNProngs) // 2 prongs vertices with one of charges set to 0: Helix : Line { std::vector forceQ{1, 1}; - LOG(info) << "Processing 2-prong Helix - Line case"; + LOG(info) << "\n\nProcessing 2-prong Helix - Line case"; o2::vertexing::DCAFitterN<2> ft; // 2 prong fitter ft.setBz(bz); ft.setPropagateToPCA(true); // After finding the vertex, propagate tracks to the DCA. This is default anyway @@ -377,17 +377,17 @@ BOOST_AUTO_TEST_CASE(DCAFitterNProngs) nfoundW++; } } - ft.print(); + // ft.print(); meanDA /= nfoundA ? nfoundA : 1; meanDAW /= nfoundAW ? nfoundAW : 1; meanDW /= nfoundW ? nfoundW : 1; LOG(info) << "Processed " << NTest << " 2-prong vertices: Helix : Line"; LOG(info) << "2-prongs with abs.dist minization: eff= " << float(nfoundA) / NTest - << " mean.dist to truth: " << meanDA << " CPU time: " << swA.CpuTime(); + << " mean.dist to truth: " << meanDA << " CPU time: " << swA.CpuTime() * 1000 << " ms"; LOG(info) << "2-prongs with abs.dist but wghPCA: eff= " << float(nfoundAW) / NTest - << " mean.dist to truth: " << meanDAW << " CPU time: " << swAW.CpuTime(); + << " mean.dist to truth: " << meanDAW << " CPU time: " << swAW.CpuTime() * 1000 << " ms"; LOG(info) << "2-prongs with wgh.dist minization: eff= " << float(nfoundW) / NTest - << " mean.dist to truth: " << meanDW << " CPU time: " << swW.CpuTime(); + << " mean.dist to truth: " << meanDW << " CPU time: " << swW.CpuTime() * 1000 << " ms"; BOOST_CHECK(nfoundA > 0.99 * NTest); BOOST_CHECK(nfoundAW > 0.99 * NTest); BOOST_CHECK(nfoundW > 0.99 * NTest); @@ -399,7 +399,7 @@ BOOST_AUTO_TEST_CASE(DCAFitterNProngs) // 2 prongs vertices with both of charges set to 0: Line : Line { std::vector forceQ{0, 0}; - LOG(info) << "Processing 2-prong Line - Line case"; + LOG(info) << "\n\nProcessing 2-prong Line - Line case"; o2::vertexing::DCAFitterN<2> ft; // 2 prong fitter ft.setBz(bz); ft.setPropagateToPCA(true); // After finding the vertex, propagate tracks to the DCA. This is default anyway @@ -454,17 +454,17 @@ BOOST_AUTO_TEST_CASE(DCAFitterNProngs) nfoundW++; } } - ft.print(); + // ft.print(); meanDA /= nfoundA ? nfoundA : 1; meanDAW /= nfoundAW ? nfoundAW : 1; meanDW /= nfoundW ? nfoundW : 1; LOG(info) << "Processed " << NTest << " 2-prong vertices: Line : Line"; LOG(info) << "2-prongs with abs.dist minization: eff= " << float(nfoundA) / NTest - << " mean.dist to truth: " << meanDA << " CPU time: " << swA.CpuTime(); + << " mean.dist to truth: " << meanDA << " CPU time: " << swA.CpuTime() * 1000 << " ms"; LOG(info) << "2-prongs with abs.dist but wghPCA: eff= " << float(nfoundAW) / NTest - << " mean.dist to truth: " << meanDAW << " CPU time: " << swAW.CpuTime(); + << " mean.dist to truth: " << meanDAW << " CPU time: " << swAW.CpuTime() * 1000 << " ms"; LOG(info) << "2-prongs with wgh.dist minization: eff= " << float(nfoundW) / NTest - << " mean.dist to truth: " << meanDW << " CPU time: " << swW.CpuTime(); + << " mean.dist to truth: " << meanDW << " CPU time: " << swW.CpuTime() * 1000 << " ms"; BOOST_CHECK(nfoundA > 0.99 * NTest); BOOST_CHECK(nfoundAW > 0.99 * NTest); BOOST_CHECK(nfoundW > 0.99 * NTest); @@ -475,8 +475,8 @@ BOOST_AUTO_TEST_CASE(DCAFitterNProngs) // 3 prongs vertices { + LOG(info) << "\n\nProcessing 3-prong vertices"; std::vector forceQ{1, 1, 1}; - o2::vertexing::DCAFitterN<3> ft; // 3 prong fitter ft.setBz(bz); ft.setPropagateToPCA(true); // After finding the vertex, propagate tracks to the DCA. This is default anyway @@ -530,17 +530,17 @@ BOOST_AUTO_TEST_CASE(DCAFitterNProngs) nfoundW++; } } - ft.print(); + // ft.print(); meanDA /= nfoundA ? nfoundA : 1; meanDAW /= nfoundAW ? nfoundAW : 1; meanDW /= nfoundW ? nfoundW : 1; LOG(info) << "Processed " << NTest << " 3-prong vertices"; LOG(info) << "3-prongs with abs.dist minization: eff= " << float(nfoundA) / NTest - << " mean.dist to truth: " << meanDA << " CPU time: " << swA.CpuTime(); + << " mean.dist to truth: " << meanDA << " CPU time: " << swA.CpuTime() * 1000 << " ms"; LOG(info) << "3-prongs with abs.dist but wghPCA: eff= " << float(nfoundAW) / NTest - << " mean.dist to truth: " << meanDAW << " CPU time: " << swAW.CpuTime(); + << " mean.dist to truth: " << meanDAW << " CPU time: " << swAW.CpuTime() * 1000 << " ms"; LOG(info) << "3-prongs with wgh.dist minization: eff= " << float(nfoundW) / NTest - << " mean.dist to truth: " << meanDW << " CPU time: " << swW.CpuTime(); + << " mean.dist to truth: " << meanDW << " CPU time: " << swW.CpuTime() * 1000 << " ms"; BOOST_CHECK(nfoundA > 0.99 * NTest); BOOST_CHECK(nfoundAW > 0.99 * NTest); BOOST_CHECK(nfoundW > 0.99 * NTest);