Alpaka Multi-depth particle flow clusterizer #48248
Conversation
…-mdpf-update4 branch)
|
cms-bot internal usage |
|
-code-checks Logs: https://cmssdt.cern.ch/SDT/code-checks/cms-sw-PR-48248/45055
Code check has found code style and quality issues which could be resolved by applying following patch(s)
|
|
assign heterogeneous |
| //SOA_COLUMN(int, adjacencyList),//for multi-depth clusterizer | ||
| //SOA_COLUMN(int, nDegree),//for multi-depth clusterizer |
There was a problem hiding this comment.
Are these commented-out lines still needed?
| //SOA_COLUMN(int, seedRHIdx), | ||
| //SOA_COLUMN(int, rhfracSize), | ||
| //SOA_COLUMN(int, rhfracOffset), |
There was a problem hiding this comment.
Are these commented0out lines still needed?
| //SOA_COLUMN(int, mdpf_component), //list of component contents (always start with component root id) | ||
| //SOA_COLUMN(int, mdpf_componentIndex), | ||
| //SOA_COLUMN(int, mdpf_componentEnergy), | ||
| //SOA_SCALAR(int, mdpf_nTopos), |
There was a problem hiding this comment.
Are these commented-out lines still needed?
There was a problem hiding this comment.
In principle the contents of this file look generic, and would be better placed in HeterogeneousCore/AlpakaInterface/interface. On the other hand, given that the AMD implementation is simple (I guess relying on __builtin_amdgcn_wave_barrier() providing stronger synchronization than minimally needed), maybe in the interest of time this place would be good-enough for now?
There was a problem hiding this comment.
Although because all files presently added to RecoParticleFlow/PFClusterProducer/interface/alpaka/ are used only in RecoParticleFlow/PFClusterProducer/plugins/alpaka, maybe all those files should be moved to plugins/alpaka directory.
There was a problem hiding this comment.
It's too early to move this to HeterogeneousCore/AlpakaInterface/interface, but I agree it makes sense to relocate it to the plugin for now.
|
|
||
| #include "HeterogeneousCore/AlpakaInterface/interface/config.h" | ||
|
|
||
| namespace ALPAKA_ACCELERATOR_NAMESPACE { |
There was a problem hiding this comment.
Given that all functions are templated over TAcc, they could (ideally should) be placed outside of ALPAKA_ACCELERATOR_NAMESPACE.
| std::optional<reco::PFRecHitFractionDeviceCollection> outPFRHFractions; | ||
| std::optional<reco::PFClusterDeviceCollection> outPFClusters; |
There was a problem hiding this comment.
The optional could be avoided with
cms_uint32_t const nClusters = event.get(inputPFClustersNum_Token_);
cms_uint32_t const nRH = event.get(inputPFRecHitNum_Token_);
outPFClusters.emplace(nClusters, event.queue());
outPFRHFractions.emplace(nRH, event.queue());
if (nClusters > 0) {
clusterizer_->apply(event.queue(), outPFClusters, outPFRHFractions, pfClusters, pfRecHitFractions, pfRecHits);
}| enum class ClusterParamKind { DEPTH, ENERGY, ETA, PHI, ETA_RMS2, PHI_RMS2, INVALID_KIND }; | ||
|
|
||
| class ClusterParam { | ||
| protected: |
There was a problem hiding this comment.
Why protected instead of private? Is this class intended to be inherited from? (I didn't quickly see any inheritance)
(same question for all other use of protected)
| template <typename TAcc, typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>> | ||
| ALPAKA_FN_ACC void operator()(TAcc const& acc, |
There was a problem hiding this comment.
Can this kernel operate on any number of dimensions? If it is restricted to one (Acc1D that is used now), how about
| template <typename TAcc, typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>> | |
| ALPAKA_FN_ACC void operator()(TAcc const& acc, | |
| ALPAKA_FN_ACC void operator()(Acc1D const& acc, |
(same comment for other kernels)
| } | ||
|
|
||
| template <ClusterParamKind kind = ClusterParamKind::INVALID_KIND> | ||
| inline constexpr auto Get() const { |
There was a problem hiding this comment.
What is the benefit of this approach compared to member functions along
constexpr auto depth() const { return depth_; }?
|
|
||
| enum class ClusterParamKind { DEPTH, ENERGY, ETA, PHI, ETA_RMS2, PHI_RMS2, INVALID_KIND }; | ||
|
|
||
| class ClusterParam { |
There was a problem hiding this comment.
Many names in this file sound generic enough for me to worry about possible symbol collisions. I'd suggest either more specific names, or a namespace specific to PF.
|
-code-checks Logs: https://cmssdt.cern.ch/SDT/code-checks/cms-sw-PR-48248/45071
Code check has found code style and quality issues which could be resolved by applying following patch(s)
|
|
-code-checks Logs: https://cmssdt.cern.ch/SDT/code-checks/cms-sw-PR-48248/45077
Code check has found code style and quality issues which could be resolved by applying following patch(s)
|
…gorithm implementation, removed misleading comments, and renamed class names to avoid potential name collisions.
|
-code-checks Logs: https://cmssdt.cern.ch/SDT/code-checks/cms-sw-PR-48248/45087 ERROR: Build errors found during clang-tidy run. |
|
-code-checks Logs: https://cmssdt.cern.ch/SDT/code-checks/cms-sw-PR-48248/45089 ERROR: Build errors found during clang-tidy run. |
|
-code-checks Logs: https://cmssdt.cern.ch/SDT/code-checks/cms-sw-PR-48248/45138
Code check has found code style and quality issues which could be resolved by applying following patch(s)
|
|
-code-checks Logs: https://cmssdt.cern.ch/SDT/code-checks/cms-sw-PR-48248/45169
Code check has found code style and quality issues which could be resolved by applying following patch(s)
|
PR description:
This PR introduces an Alpaka-based implementation of the multi-depth particle flow clustering algorithm, which makes extensive use of warp-level operations. Due to the inherently non-local nature of the algorithm, we implemented custom masked warp intrinsics currently not available in the Alpaka library. Additionally, we utilized the extended ECL-CC (connected component detection) graph algorithm to support this use case. This extension involves reorganizing SoA data objects into a CSR (Compressed Sparse Row) format during both pre-processing and post-processing stages to enable efficient graph traversal.
PR validation:
The implementation expects input data to be provided in device-resident SoA format.