Log Report for PR reviews

[AdrianoDee]

This issue keeps track of the comment addressed for the PRs:

• Porting Pixel Tracks to Alpaka [Not to Merge]  cms-sw/cmssw#41117 (superseded by the following)
• Pixel Alpaka Migration: Modification to Existing Modules [I] cms-sw/cmssw#41282
• Pixel Alpaka Migration: AlpakaInterface Updates [II] cms-sw/cmssw#41284
• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285
• Pixel Alpaka Migration: Tracking [IV] cms-sw/cmssw#41286
• Pixel Alpaka Migration: Vertexing [V] cms-sw/cmssw#41287
• Pixel Alpaka Migration: DQM [VI] cms-sw/cmssw#41288

---

23 March 2023

• Porting Pixel Tracks to Alpaka [Not to Merge]  cms-sw/cmssw#41117 (comment) by @tvami
  

  cmssw/CalibTracker/SiPixelESProducers/interface/SiPixelESTestData.h

  Line 1 in 58b4bba

  #ifndef CalibTracker_SiPixelES_interface_PixelESTestData_h

  • Change with:

   #ifndef CalibTracker_SiPixelESProducers_SiPixelESTestData_h

• Porting Pixel Tracks to Alpaka [Not to Merge]  cms-sw/cmssw#41117 (comment) by @makortel
  

  cmssw/CalibTracker/Records/BuildFile.xml

  Lines 6 to 8 in 58b4bba

  	<use name="HeterogeneousCore/AlpakaCore"/>
  	<use name="HeterogeneousCore/AlpakaInterface"/>
  	<flags ALPAKA_BACKENDS="1"/>

  • Suggestion: Dependence on Alpaka is not needed for Record definitions

• Porting Pixel Tracks to Alpaka [Not to Merge]  cms-sw/cmssw#41117 (comment) by @makortel
  

  cmssw/CalibTracker/SiPixelESProducers/interface/SiPixelGainCalibrationForHLTHost.h

  Lines 4 to 5 in 58b4bba

  	#include <cstdint>
  	#include <alpaka/alpaka.hpp>

  • Suggestion: These includes seem to unnecessary

• Porting Pixel Tracks to Alpaka [Not to Merge]  cms-sw/cmssw#41117 (comment) by @makortel
  

  cmssw/CalibTracker/SiPixelESProducers/interface/SiPixelGainCalibrationForHLTHost.h

  Line 7 in 58b4bba

  #include "SiPixelGainCalibrationForHLTLayout.h"

  
  Our convention for files in interface is to use full path to headers.

• Porting Pixel Tracks to Alpaka [Not to Merge]  cms-sw/cmssw#41117 (comment) by @makortel
  

  cmssw/CalibTracker/SiPixelESProducers/interface/SiPixelGainCalibrationForHLTLayout.h

  Line 19 in 58b4bba

  GENERATE_SOA_LAYOUT(SiPixelGainCalibrationForHLTLayout,

  
  The SoA layout and the corresponding PortableCollections correspond to data formats, and I think they should be placed as such. On a cursory look CondFormats/SiPixelObjects would look like a plausible candidate (even if none of these are put in the CondDB).

• Porting Pixel Tracks to Alpaka [Not to Merge]  cms-sw/cmssw#41117 (comment) by @makortel
  

  cmssw/CalibTracker/Records/interface/SiPixelESTestRecords.h

  Line 8 in 58b4bba

  class PixelESTestRecord

  
  I assume all this kind of test code gets removed in the end.

• Porting Pixel Tracks to Alpaka [Not to Merge]  cms-sw/cmssw#41117 (comment) by @makortel
  

  cmssw/CalibTracker/SiPixelESProducers/plugins/alpaka/SiPixelCablingSoAESProducer.cc

  Line 28 in 58b4bba

  template struct cms::alpakatools::CopyToDevice<SiPixelMappingHost>; //needed for the method to not be incomplete

  
  I'm confused why this is needed.

• Porting Pixel Tracks to Alpaka [Not to Merge]  cms-sw/cmssw#41117 (comment) by @makortel
  

  cmssw/CalibTracker/SiPixelESProducers/plugins/alpaka/SiPixelCablingSoAESProducer.cc

  Lines 38 to 39 in 58b4bba

  	auto const& component = iConfig.getParameter<std::string>("ComponentName");
  	auto cc = setWhatProduced(this, component);

  
  In principle this ComponentName is redundant and the ESProducer could just rely on the framework-provided appendToDataLabel, but I don't know if the ship sailed already.

• Porting Pixel Tracks to Alpaka [Not to Merge]  cms-sw/cmssw#41117 (comment) by @makortel
  

  cmssw/DQM/SiPixelHeterogeneous/plugins/alpaka/SiPixelCompareTrackSoA.cc

  Line 29 in 58b4bba

  // cf https://github.com/cms-sw/cmssw/blob/master/Validation/RecoTrack/src/MTVHistoProducerAlgoForTracker.cc

  
  Maybe it would be a time to provide these from some common area?

• Porting Pixel Tracks to Alpaka [Not to Merge]  cms-sw/cmssw#41117 (comment) by @makortel
  

  cmssw/DQM/SiPixelHeterogeneous/plugins/alpaka/SiPixelCompareTrackSoA.cc

  Line 66 in 58b4bba

  class SiPixelCompareTrackSoA : public DQMEDAnalyzer {

  
  The DQMEDAnalyzer should, in the leading order, be independent of Alpaka and read the host product. Did that not work for some reason?

• Porting Pixel Tracks to Alpaka [Not to Merge]  cms-sw/cmssw#41117 (comment) by @makortel
  

  cmssw/DataFormats/GeometrySurface/interface/SOARotation.h

  Line 15 in 58b4bba

  inline SOARotation() {}

  
  Why is constexpr removed?

• Porting Pixel Tracks to Alpaka [Not to Merge]  cms-sw/cmssw#41117 (comment) by @makortel
  

  cmssw/DataFormats/PixelCPEFastParams/src/PixelCPEFastParams.cc

  Line 9 in 58b4bba

  TYPELOOKUP_DATA_REG(PixelCPEFastParamsPhase2);

  
  Do these EventSetup data products really need to be in DataFormats? DataFormats is primarily for Event data products (although I see two EventSetup data products currently defined in DataFormats, IIRC both have some dependence-related rationale)

• Porting Pixel Tracks to Alpaka [Not to Merge]  cms-sw/cmssw#41117 (comment) by @makortel
  

  cmssw/DataFormats/Portable/interface/HostProduct.h

  Line 10 in 58b4bba

  class HostProduct {

  
  Is this a temporary measure to help the transition, or intended for long term?

• Porting Pixel Tracks to Alpaka [Not to Merge]  cms-sw/cmssw#41117 (comment) by @makortel
  

  cmssw/DataFormats/SiPixelClusterSoA/src/alpaka/classes_serial_def.xml

  Line 2 in 58b4bba

  <class name="PortableHostCollection<SiPixelClustersLayout>"/>

  
  Please move the host collection dictionary declarations from alpaka/classes_serial{.h,_def.xml} to classes{.h,_def.xml} (the placement convention was updated).

• Porting Pixel Tracks to Alpaka [Not to Merge]  cms-sw/cmssw#41117 (comment) by @makortel
  

  cmssw/RecoTauTag/HLTProducers/src/L2TauTagNNProducer.cc

  Line 631 in 58b4bba

  ALPAKA_ACCELERATOR_NAMESPACE::TracksUtilities<pixelTopology::Phase1>::copyToDense(

  
  Non-alpaka code must stay independent of ALPAKA_ACCELERATOR_NAMESPACE.

---

28 March 2023

• Porting Pixel Tracks to Alpaka [Not to Merge]  cms-sw/cmssw#41117 (comment) by @fwyzard
  

  cmssw/DataFormats/TrackingRecHitSoA/interface/alpaka/TrackingRecHitSoADevice.h

  Line 42 in 58b4bba

  const auto device = cms::alpakatools::devices<Platform>()[0];

  
  This is wrong. Please use the device associated to the queue.

• Porting Pixel Tracks to Alpaka [Not to Merge]  cms-sw/cmssw#41117 (comment) by @fwyzard
  

  cmssw/DataFormats/Track/interface/alpaka/PixelTrackUtilities.h

  Line 22 in 58b4bba

  constexpr ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE static float phi(const TrackSoAConstView &tracks, int32_t i) {

  • Change with:

       ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE static constexpr float phi(const TrackSoAConstView &tracks, int32_t i) {

• Porting Pixel Tracks to Alpaka [Not to Merge]  cms-sw/cmssw#41117 (comment) by @fwyzard
  

  cmssw/DataFormats/Track/interface/alpaka/PixelTrackUtilities.h

  Line 26 in 58b4bba

  constexpr ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE static float tip(const TrackSoAConstView &tracks, int32_t i) {

  • Change with:

       ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE static constexpr float tip(const TrackSoAConstView &tracks, int32_t i) {

• Porting Pixel Tracks to Alpaka [Not to Merge]  cms-sw/cmssw#41117 (comment) by @fwyzard
  

  cmssw/DataFormats/Track/interface/alpaka/PixelTrackUtilities.h

  Line 30 in 58b4bba

  constexpr ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE static float zip(const TrackSoAConstView &tracks, int32_t i) {

  • Change with:

       ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE static constexpr float zip(const TrackSoAConstView &tracks, int32_t i) {

• Porting Pixel Tracks to Alpaka [Not to Merge]  cms-sw/cmssw#41117 (comment) by @fwyzard
  

  cmssw/DataFormats/Track/interface/alpaka/PixelTrackUtilities.h

  Line 34 in 58b4bba

  constexpr ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE static bool isTriplet(const TrackSoAConstView &tracks, int i) {

  • Change with:

       ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE static constexpr bool isTriplet(const TrackSoAConstView &tracks, int i) {

• Porting Pixel Tracks to Alpaka [Not to Merge]  cms-sw/cmssw#41117 (comment) by @fwyzard
  

  cmssw/DataFormats/Track/interface/alpaka/PixelTrackUtilities.h

  Line 39 in 58b4bba

  constexpr ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE static void copyFromCircle(

  • Change with:

       ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE static constexpr void copyFromCircle(

• Porting Pixel Tracks to Alpaka [Not to Merge]  cms-sw/cmssw#41117 (comment) by @fwyzard
  

  cmssw/DataFormats/Track/interface/alpaka/PixelTrackUtilities.h

  Line 60 in 58b4bba

  constexpr ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE static void copyFromDense(TrackSoAView &tracks,

  • Change with:

       ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE static constexpr void copyFromDense(TrackSoAView &tracks,

• Porting Pixel Tracks to Alpaka [Not to Merge]  cms-sw/cmssw#41117 (comment) by @fwyzard
  

  cmssw/DataFormats/Track/interface/alpaka/PixelTrackUtilities.h

  Line 71 in 58b4bba

  constexpr ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE static void copyToDense(const TrackSoAConstView &tracks,

  • Change with:

       ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE static constexpr void copyToDense(const TrackSoAConstView &tracks,

• Porting Pixel Tracks to Alpaka [Not to Merge]  cms-sw/cmssw#41117 (comment) by @fwyzard
  

  cmssw/DataFormats/TrackingRecHitSoA/interface/alpaka/TrackingRecHitSoADevice.h

  Lines 41 to 42 in 58b4bba

  	const auto host = cms::alpakatools::host();
  	const auto device = cms::alpakatools::devices<Platform>()[0];

  • Change with:

         const auto& host = cms::alpakatools::host();
         const auto device = alpaka::getDev(queue);

• Porting Pixel Tracks to Alpaka [Not to Merge]  cms-sw/cmssw#41117 (comment) by @fwyzard
  

  cmssw/DataFormats/TrackingRecHitSoA/interface/alpaka/TrackingRecHitSoADevice.h

  Line 50 in 58b4bba

  alpaka::memcpy(queue, start_d, start_h);//, TrackerTraits::numberOfModules + 1);

  • Change with:

         alpaka::memcpy(queue, start_d, start_h);

• Porting Pixel Tracks to Alpaka [Not to Merge]  cms-sw/cmssw#41117 (comment) by @fwyzard
  

  cmssw/DataFormats/Track/interface/alpaka/PixelTrackUtilities.h

  Line 18 in 58b4bba

  ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE static float charge(const TrackSoAConstView &tracks, int32_t i) {

  
  Is this function actually used ?

  In principle, std::copysign() should not work in e.g. CUDA device code, because it's not constexpr.

• Porting Pixel Tracks to Alpaka [Not to Merge]  cms-sw/cmssw#41117 (comment) by @makortel
  

  cmssw/DataFormats/Track/interface/alpaka/PixelTrackUtilities.h

  Line 8 in 58b4bba

  namespace ALPAKA_ACCELERATOR_NAMESPACE {

  
  As far as I can tell, nothing in this file needs to be in ALPAKA_ACCELERATOR_NAMESPACE.

---

31 March 2023

• Porting Pixel Tracks to Alpaka [Not to Merge]  cms-sw/cmssw#41117 (comment) by @makortel
  

  cmssw/CalibTracker/SiPixelESProducers/interface/SiPixelMappingSoAESRecord.h

  Line 7 in 58b4bba

  class SiPixelMappingSoAESRecord : public edm::eventsetup::EventSetupRecordImplementation<SiPixelMappingSoAESRecord> {};

  
  This Record seems to be unused

• Porting Pixel Tracks to Alpaka [Not to Merge]  cms-sw/cmssw#41117 (comment) by @makortel
  

  cmssw/CalibTracker/SiPixelESProducers/interface/alpaka/SiPixelGainCalibrationForHLTUtilities.h

  Lines 13 to 19 in 58b4bba

  	ALPAKA_FN_HOST_ACC ALPAKA_FN_ACC ALPAKA_FN_INLINE static std::pair<float, float> getPedAndGain(
  	uint32_t moduleInd,
  	int col,
  	int row,
  	bool& isDeadColumn,
  	bool& isNoisyColumn,
  	const SiPixelGainCalibrationForHLTSoAConstView& view) {

  
  This function looks like in the AoS model it would be a member function of the struct. I suggest we treat these kind of functions for SoAs such that
  • they are defined in the same namespace (and file) as the Layout and View are defined (i.e. where the GENERATE_SOA_LAYOUT() macro is called)
  • the View object is the first parameter (mimicking the implicit this argument of actual member functions)

  The function call would then look like

   getPedAndGain(view, ...);

  which would not be very different from

   view.getPedAndGain(...);

  (by the way, I don't see anything in this file that would require the use of ALPAKA_ACCELERATOR_NAMESPACE)

• Porting Pixel Tracks to Alpaka [Not to Merge]  cms-sw/cmssw#41117 (comment) by @makortel
  

  cmssw/CalibTracker/SiPixelESProducers/interface/alpaka/SiPixelMappingUtilities.h

  Lines 27 to 28 in 58b4bba

  	std::vector<unsigned int> const& fedIds = cablingMap->fedIds();
  	std::unique_ptr<SiPixelFedCablingTree> const& cabling = cablingMap->cablingTree();

  
  How about passing these two from the caller to avoid constructing them for every event, since the caller (SiPixelRawToCluster) already caches them by IOVs?

• Porting Pixel Tracks to Alpaka [Not to Merge]  cms-sw/cmssw#41117 (comment) by @makortel
  

  cmssw/CalibTracker/SiPixelESProducers/interface/alpaka/SiPixelMappingUtilities.h

  Lines 30 to 31 in 58b4bba

  	unsigned int startFed = *(fedIds.begin());
  	unsigned int endFed = *(fedIds.end() - 1);

  • Change with:

         unsigned int startFed = fedIds.front();
         unsigned int endFed = fedIds.back();

  would be more clear

• Porting Pixel Tracks to Alpaka [Not to Merge]  cms-sw/cmssw#41117 (comment) by @makortel
  

  cmssw/CalibTracker/SiPixelESProducers/interface/alpaka/SiPixelMappingUtilities.h

  Line 25 in 58b4bba

  auto modToUnpHost = cms::alpakatools::make_device_buffer<unsigned char[]>(queue, pixelgpudetails::MAX_SIZE);

  • Change with:
    Should this be

         auto modToUnpHost = cms::alpakatools::make_host_buffer<unsigned char[]>(queue, pixelgpudetails::MAX_SIZE);

  ?

---

03 April 2023

• Porting Pixel Tracks to Alpaka [Not to Merge]  cms-sw/cmssw#41117 (comment) by @makortel
  

  cmssw/CalibTracker/Records/interface/SiPixelMappingSoARecord.h

  Line 7 in 99f9438

  class SiPixelMappingSoARecord : public edm::eventsetup::EventSetupRecordImplementation<SiPixelMappingSoARecord> {};

  
  I don't see any ESProducer producing the SiPixelMappingDevice or SiPixelMappingHost into this SiPixelMappingSoARecord record.

  I see SiPixelCablingSoAESProducer produces SiPixelMappingHost into CkfComponentsRecord. Is the intention to use CkfComponentsRecord (in which case this record could be removed), or something more targeted (in which case this record needs to be made to depend on SiPixelFedCablingMapRcd, SiPixelQualityRcd, and TrackerDigiGeometryRecord)?

• Porting Pixel Tracks to Alpaka [Not to Merge]  cms-sw/cmssw#41117 (comment) by @makortel
  

  cmssw/CalibTracker/SiPixelESProducers/interface/SiPixelMappingHost.h

  Line 4 in 99f9438

  #include <cstdint>

  • Suggestion: This #include does not seem to be needed

• Porting Pixel Tracks to Alpaka [Not to Merge]  cms-sw/cmssw#41117 (comment) by @makortel
  

  cmssw/CalibTracker/SiPixelESProducers/interface/SiPixelMappingHost.h

  Line 7 in 99f9438

  #include "SiPixelMappingLayout.h"

  • Change with:
    Full path

   #include "CalibTracker/SiPixelESProducers/interface/SiPixelMappingLayout.h"

• Porting Pixel Tracks to Alpaka [Not to Merge]  cms-sw/cmssw#41117 (comment) by @makortel
  

  cmssw/CalibTracker/SiPixelESProducers/interface/SiPixelMappingLayout.h

  Line 8 in 99f9438

  GENERATE_SOA_LAYOUT(SiPixelMappingLayout,

  
  Maybe this SoA should be defined in a data format package as well?

• Porting Pixel Tracks to Alpaka [Not to Merge]  cms-sw/cmssw#41117 (comment) by @makortel
  

  cmssw/CalibTracker/SiPixelESProducers/interface/SiPixelStatusAndCablingLayout.h

  Line 6 in 99f9438

  GENERATE_SOA_LAYOUT(SiPixelStatusAndCablingLayout,

  
  Move to a data format package?

• Porting Pixel Tracks to Alpaka [Not to Merge]  cms-sw/cmssw#41117 (comment) by @makortel
  

  cmssw/CondFormats/SiPixelObjects/BuildFile.xml

  Line 17 in 99f9438

  <use name="HeterogeneousCore/AlpakaCore"/>

  
  Dependence on AlpakaCore should not be needed in a data format package

• Porting Pixel Tracks to Alpaka [Not to Merge]  cms-sw/cmssw#41117 (comment) by @makortel
  

  cmssw/CondFormats/SiPixelObjects/BuildFile.xml

  Line 19 in 99f9438

  <use name="HeterogeneousCore/CUDACore"/>

  
  Dependence on CUDACore should not be needed in a data format package

• Porting Pixel Tracks to Alpaka [Not to Merge]  cms-sw/cmssw#41117 (comment) by @makortel
  

  cmssw/CondFormats/SiPixelObjects/BuildFile.xml

  Line 2 in 99f9438

  <use name="cuda"/>

  
  Why is dependence on cuda added?

• Porting Pixel Tracks to Alpaka [Not to Merge]  cms-sw/cmssw#41117 (comment) by @makortel
  

  cmssw/CondFormats/SiPixelObjects/interface/SiPixelGainCalibrationForHLTHost.h

  Line 2 in 99f9438
  • Suggestion:

• Porting Pixel Tracks to Alpaka [Not to Merge]  cms-sw/cmssw#41117 (comment) by @makortel
  

  cmssw/CondFormats/SiPixelObjects/interface/alpaka/SiPixelGainCalibrationForHLTDevice.h

  Line 4 in 99f9438

  #include <cstdint>

  • Suggestion:

• Porting Pixel Tracks to Alpaka [Not to Merge]  cms-sw/cmssw#41117 (comment) by @makortel
  

  cmssw/CondFormats/SiPixelObjects/interface/alpaka/SiPixelGainCalibrationForHLTUtilities.h

  Line 7 in 99f9438

  #include "DataFormats/Portable/interface/alpaka/PortableCollection.h"

  • Suggestion: Does not seem to be needed

• Porting Pixel Tracks to Alpaka [Not to Merge]  cms-sw/cmssw#41117 (comment) by @makortel
  

  cmssw/CondFormats/SiPixelObjects/interface/alpaka/SiPixelGainCalibrationForHLTUtilities.h

  Line 1 in 99f9438
  • Suggestion:

• Porting Pixel Tracks to Alpaka [Not to Merge]  cms-sw/cmssw#41117 (comment) by @makortel
  

  cmssw/CondFormats/SiPixelObjects/interface/alpaka/SiPixelGainCalibrationForHLTUtilities.h

  Line 1 in 99f9438
  
  Now that this file does not use ALPAKA_ACCELERATOR_NAMESPACE, it should be moved one level up to CondFormats/SiPixelObjects/interface/.

  Although I would still move the getPedAndGain() as a free function to CondFormats/SiPixelObjects/interface/SiPixelGainCalibrationForHLTLayout.h.

• Porting Pixel Tracks to Alpaka [Not to Merge]  cms-sw/cmssw#41117 (comment) by @makortel
  

  cmssw/DQM/SiPixelHeterogeneous/plugins/BuildFile.xml

  Line 1 in 99f9438

  <use name="alpaka"/>

  
  Why explicit dependence on alpaka?

• Porting Pixel Tracks to Alpaka [Not to Merge]  cms-sw/cmssw#41117 (comment) by @makortel
  

  cmssw/DataFormats/BeamSpotAlpaka/interface/alpaka/BeamSpotAlpaka.h

  Line 29 in 99f9438

  cms::alpakatools::device_buffer<Device, BeamSpotPOD> const& ptr() const { return *data_d_; }

  
  IIUC the device_buffer maps underneath to alpaka::Buf, so this function would allow non-const access e.g. via

   cms::alpakatools::device_buffer<Device, BeamSpotPOD> tmp = beamspot.ptr();
   // tmp gives non-const access to the BesmSpotPOD owned by beamspot

  that violates data format class requirements.

• Porting Pixel Tracks to Alpaka [Not to Merge]  cms-sw/cmssw#41117 (comment) by @makortel
  

  cmssw/DataFormats/BeamSpotAlpaka/src/alpaka/classes_serial.h

  Line 1 in 99f9438

  #ifndef DataFormats_BeamSpotAlpaka_alpaka_classes_serial_h

  
  Please move the contents to DataFormats/BeamSpotAlpaka/src/classes.h

• Porting Pixel Tracks to Alpaka [Not to Merge]  cms-sw/cmssw#41117 (comment) by @makortel
  

  cmssw/DataFormats/BeamSpotAlpaka/src/alpaka/classes_serial_def.xml

  Line 1 in 99f9438

  <lcgdict>

  
  Please move the contents to DataFormats/BeamSpotAlpaka/src/classes_def.xml

• Porting Pixel Tracks to Alpaka [Not to Merge]  cms-sw/cmssw#41117 (comment) by @makortel
  

  cmssw/DataFormats/BeamSpotAlpaka/BuildFile.xml

  Line 5 in 99f9438

  <flags ALPAKA_BACKENDS="serial"/>

  • Change with:
    Should this be

   <flags ALPAKA_BACKENDS="cuda rocm"/>

  ?

• Porting Pixel Tracks to Alpaka [Not to Merge]  cms-sw/cmssw#41117 (comment) by @makortel
  

  cmssw/DataFormats/PixelCPEFastParams/BuildFile.xml

  Lines 2 to 3 in 99f9438

  	<use name="FWCore/Framework"/>
  	<use name="HeterogeneousCore/AlpakaCore"/>

  
  Data format package must not depend on Framework or AlpakaCore.

• Porting Pixel Tracks to Alpaka [Not to Merge]  cms-sw/cmssw#41117 (comment) by @makortel
  

  cmssw/DataFormats/PixelCPEFastParams/BuildFile.xml

  Line 6 in 99f9438

  <use name="RecoLocalTracker/SiPixelRecHits"/>

  
  Data format package must not depend on RecoLocalTracker

• Porting Pixel Tracks to Alpaka [Not to Merge]  cms-sw/cmssw#41117 (comment) by @makortel
  

  cmssw/DataFormats/PixelCPEFastParams/BuildFile.xml

  Line 7 in 99f9438

  <flags ALPAKA_BACKENDS="1"/>

  • Change with:
    Should probably be

   <flags ALPAKA_BACKENDS="cuda rocm"/>

• Porting Pixel Tracks to Alpaka [Not to Merge]  cms-sw/cmssw#41117 (comment) by @makortel
  

  cmssw/DataFormats/SiPixelClusterSoA/BuildFile.xml

  Line 6 in 99f9438

  <flags ALPAKA_BACKENDS="1"/>

  • Change with:

   <flags ALPAKA_BACKENDS="cuda rocm"/>

• Porting Pixel Tracks to Alpaka [Not to Merge]  cms-sw/cmssw#41117 (comment) by @makortel
  

  cmssw/CondFormats/SiPixelObjects/BuildFile.xml

  Line 20 in 99f9438

  <flags ALPAKA_BACKENDS="1"/>

  • Change with:

   <flags ALPAKA_BACKENDS="cuda rocm"/>

• Porting Pixel Tracks to Alpaka [Not to Merge]  cms-sw/cmssw#41117 (comment) by @makortel
  

  cmssw/DataFormats/SiPixelClusterSoA/interface/alpaka/SiPixelClustersDevice.h

  Line 16 in 99f9438

  class SiPixelClustersDevice : public PortableCollection<SiPixelClustersLayout<>> {

  
  The problem here is that alpaka_serial_sync::SiPixelClustersDevice is a different type than theSiPixelClustersHost (I'll clarify this requirement in the README).

  The PortableCollection itself is a backend-dependent type alias to PortableCollectionHost and PortableCollectionDevice
  https://github.com/cms-sw/cmssw/blob/c4038b4c091680c64a4b4b8876084996ac59d0a8/DataFormats/Portable/interface/alpaka/PortableCollection.h#L17-L30

  A possible fix could be

   #ifdef ALPAKA_ACC_CPU_B_SEQ_T_SEQ_ENABLED
     using SiPixelClustersDevice = SiPixelClustersHost;
   #else
     class SiPixelClustersDevice : public PortableCollection<SiPixelClustersLayout<>> {
     ...
   #endif

  (although if we continue this with #ifdef, I think we should consider defining our own for "any host backend". macro)

  An alternative could be

     class SiPixelClustersDeviceImpl : public PortableCollection<SiPixelClustersLayout<>> { ... };
     using SiPixelClustersDevice = std::conditional_t<std::is_same_v<Platform, alpaka::PltfCpu>, SiPixelClustersHost, SiPixelClustersDeviceImpl>;

  although in this case I think we should consider abstracting the std::conditional_t<std::is_same_v<Platform, alpaka::PltfCpu>, HostType, DeviceType> part (since the use case seems to be more widespread than PortableCollection).

• Porting Pixel Tracks to Alpaka [Not to Merge]  cms-sw/cmssw#41117 (comment) by @makortel
  

  cmssw/DataFormats/SiPixelClusterSoA/src/alpaka/classes_rocm_def.xml

  Line 3 in 99f9438

  <class name="edm::Wrapper<alpaka_rocm_async::SiPixelClustersDevice>" persistent="false"/>

  • Suggestion:

• Porting Pixel Tracks to Alpaka [Not to Merge]  cms-sw/cmssw#41117 (comment) by @makortel
  

  cmssw/DataFormats/SiPixelClusterSoA/src/alpaka/classes_cuda_def.xml

  Line 3 in 99f9438

  <class name="edm::Wrapper<alpaka_cuda_async::SiPixelClustersDevice>" persistent="false"/>

  • Suggestion:

• Porting Pixel Tracks to Alpaka [Not to Merge]  cms-sw/cmssw#41117 (comment) by @makortel
  

  cmssw/DataFormats/SiPixelClusterSoA/src/alpaka/classes_serial_def.xml

  Line 1 in 99f9438

  <lcgdict>

  
  After the alpaka_serial_sync::SiPixelClustersDevice gets set up properly, this file (and classes_serial.h) should not be needed anymore.

• Porting Pixel Tracks to Alpaka [Not to Merge]  cms-sw/cmssw#41117 (comment) by @makortel
  

  cmssw/DataFormats/SiPixelClusterSoA/test/BuildFile.xml

  Line 3 in 99f9438

  <use name="HeterogeneousCore/AlpakaInterface"/>

  
  Needs also

   <use name="DataFormats/SiPixelClusterSoa" source_only="1"/>

---

04 April 2023

• Porting Pixel Tracks to Alpaka [Not to Merge]  cms-sw/cmssw#41117 (comment) by @fwyzard
  

  cmssw/HeterogeneousCore/AlpakaInterface/test/alpaka/testVec.cc

  Line 7 in 99f9438

  // #include "HeterogeneousCore/AlpakaInterface/interface/vec.h"

  
  if the include is not needed, could you remove it ?

• Porting Pixel Tracks to Alpaka [Not to Merge]  cms-sw/cmssw#41117 (comment) by @fwyzard
  

  cmssw/HeterogeneousCore/AlpakaInterface/interface/workdivision.h

  Line 10 in 99f9438

  // #include "HeterogeneousCore/AlpakaInterface/interface/vec.h"

  
  needed ?

---

05 April 2023

• Pixel Alpaka Migration: Modification to Existing Modules [I] cms-sw/cmssw#41282 (comment) by @slava77
  

  cmssw/RecoTracker/PixelSeeding/plugins/CAStructures.h

  Line 2 in cf618e5

  #define RecoPixelVertexing_PixelTriplets_CAStructures_h

  
  perhaps RecoTracker_PixelSeeding ?

• Pixel Alpaka Migration: AlpakaInterface Updates [II] cms-sw/cmssw#41284 (comment) by @makortel
  

  cmssw/HeterogeneousCore/AlpakaInterface/interface/workdivision.h

  Line 10 in 0a63573

  // #include "HeterogeneousCore/AlpakaInterface/interface/vec.h"

  • Suggestion:

• Pixel Alpaka Migration: AlpakaInterface Updates [II] cms-sw/cmssw#41284 (comment) by @makortel
  

  cmssw/HeterogeneousCore/AlpakaInterface/interface/workdivision.h

  Lines 306 to 325 in 0a63573

  	/*
  	* Computes the range of the elements indexes, local to the block.
  	* Truncated by the max number of elements of interest.
  	*/
  	// template <typename TAcc>
  	// ALPAKA_FN_ACC std::pair<Idx, Idx> element_index_range_in_block_truncated(const TAcc& acc,
  	// const Idx maxNumberOfElements,
  	// const Idx elementIdxShift,
  	// const unsigned int dimIndex = 0u) {
  	// // Check dimension
  	// //static_assert(alpaka::Dim<TAcc>::value == Dim1D::value,
  	// // "Accelerator and maxNumberOfElements need to have same dimension.");
  	// auto [firstElementIdxLocal, endElementIdxLocal] = element_index_range_in_block(acc, elementIdxShift, dimIndex);
  	// // Truncate
  	// endElementIdxLocal = std::min(endElementIdxLocal, maxNumberOfElements);
  	// // Return element indexes, shifted by elementIdxShift, and truncated by maxNumberOfElements.
  	// return {firstElementIdxLocal, endElementIdxLocal};
  	// }

  
  Is this commented code still useful?

• Pixel Alpaka Migration: AlpakaInterface Updates [II] cms-sw/cmssw#41284 (comment) by @makortel
  

  cmssw/HeterogeneousCore/AlpakaInterface/test/alpaka/testVec.cc

  Line 7 in 0a63573

  // #include "HeterogeneousCore/AlpakaInterface/interface/vec.h"

  • Suggestion:

• Pixel Alpaka Migration: AlpakaInterface Updates [II] cms-sw/cmssw#41284 (comment) by @makortel
  

  cmssw/HeterogeneousCore/AlpakaUtilities/interface/AtomicPairCounter.h

  Line 1 in 0a63573

  #ifndef AlpakaCore_AtomicPairCounter_h

  
  In the hackathon with @fwyzard and @rovere we concluded that the kind of classes in AlpakaUtilities here would be ok to be plated in AlpakaInterface (to avoid adding new package), for now at least.

• Pixel Alpaka Migration: AlpakaInterface Updates [II] cms-sw/cmssw#41284 (comment) by @makortel
  

  cmssw/HeterogeneousCore/AlpakaUtilities/interface/AtomicPairCounter.h

  Line 1 in 0a63573

  #ifndef AlpakaCore_AtomicPairCounter_h

  
  Include guard should reflect the location of the file (in all of the headers)

• Pixel Alpaka Migration: AlpakaInterface Updates [II] cms-sw/cmssw#41284 (comment) by @makortel
  

  cmssw/HeterogeneousCore/AlpakaUtilities/interface/radixSort.h

  Line 248 in 0a63573

  /* Not needed

  
  Can it be removed?

• Pixel Alpaka Migration: AlpakaInterface Updates [II] cms-sw/cmssw#41284 (comment) by @makortel
  

  cmssw/HeterogeneousCore/AlpakaUtilities/interface/demangle.h

  Line 6 in 0a63573

  namespace edm {

  
  Put the inline variable to some other namespace than edm. We have also other demanglers within CMSSW, e.g.
  https://github.com/cms-sw/cmssw/blob/107a5fed908622757511062f46961c0f7552fba8/FWCore/Utilities/interface/TypeDemangler.h#L6-L7
  (although that particular function is geared to be compatible with the convention of reflex)

• Pixel Alpaka Migration: AlpakaInterface Updates [II] cms-sw/cmssw#41284 (comment) by @makortel
  

  cmssw/HeterogeneousCore/AlpakaUtilities/interface/initialise.h

  Line 9 in 0a63573

  void initialise();

  
  The implementation file of this function is missing.

[AdrianoDee]

---

19 April 2023

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @fwyzard
  

  cmssw/CalibTracker/SiPixelESProducers/src/ES_SiPixelMappingSoA.cc

  Line 1 in 37724f7

  #include "CondFormats/SiPixelObjects/interface/SiPixelMappingHost.h"

  
  Should this file be called ES_SiPixelMappingHost.cc ?

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @fwyzard
  

  cmssw/CalibTracker/SiPixelESProducers/src/alpaka/ES_SiPixelGainCalibrationForHLTSoA.cc

  Line 1 in 37724f7

  #include "CondFormats/SiPixelObjects/interface/alpaka/SiPixelGainCalibrationForHLTDevice.h"

  
  ES_SiPixelGainCalibrationForHLTDevice.cc ?

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @fwyzard
  

  cmssw/CalibTracker/SiPixelESProducers/src/alpaka/ES_SiPixelMappingSoA.cc

  Line 1 in 37724f7

  #include "CondFormats/SiPixelObjects/interface/alpaka/SiPixelMappingDevice.h"

  
  ES_SiPixelMappingDevice.cc ?

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @fwyzard
  

  cmssw/CalibTracker/SiPixelESProducers/src/ES_SiPixelGainCalibrationForHLTSoA.cc

  Line 1 in 37724f7

  #include "CondFormats/SiPixelObjects/interface/SiPixelGainCalibrationForHLTHost.h"

  
  ES_SiPixelGainCalibrationForHLTHost.cc ?

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @fwyzard
  

  cmssw/CondFormats/SiPixelObjects/interface/SiPixelGainCalibrationForHLTHost.h

  Line 7 in 37724f7

  using SiPixelGainCalibrationForHLTHost = PortableHostCollection<SiPixelGainCalibrationForHLTLayout<>>;

  • Change with:

   using SiPixelGainCalibrationForHLTHost = PortableHostCollection<SiPixelGainCalibrationForHLTSoA>;

  In general, I would use the ...SoA alias everywhere instead of the ...Layout<> template type.

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @fwyzard
  

  cmssw/CondFormats/SiPixelObjects/interface/SiPixelGainCalibrationForHLTLayout.h

  Line 37 in 37724f7

  using SiPixelGainCalibrationForHLTSoAView = SiPixelGainCalibrationForHLTLayout<>::View;

  • Change with:
    would this work ?

   using SiPixelGainCalibrationForHLTSoAView = SiPixelGainCalibrationForHLTSoA::View;

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @fwyzard
  

  cmssw/CondFormats/SiPixelObjects/interface/SiPixelMappingLayout.h

  Line 20 in 37724f7

  using SiPixelMappingLayoutSoA = SiPixelMappingLayout<>;

  • Change with:

   using SiPixelMappingSoA = SiPixelMappingLayout<>;

  ?

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @fwyzard
  

  cmssw/CondFormats/SiPixelObjects/interface/SiPixelMappingLayout.h

  Line 21 in 37724f7

  using SiPixelMappingLayoutSoAView = SiPixelMappingLayout<>::View;

  • Change with:

   using SiPixelMappingSoAView = SiPixelMappingSoA::View;

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @fwyzard
  

  cmssw/CondFormats/SiPixelObjects/interface/SiPixelMappingLayout.h

  Line 22 in 37724f7

  using SiPixelMappingLayoutSoAConstView = SiPixelMappingLayout<>::ConstView;

  • Change with:

   using SiPixelMappingSoAConstView = SiPixelMappingSoA::ConstView;

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @fwyzard
  

  cmssw/CondFormats/SiPixelObjects/interface/alpaka/SiPixelMappingDevice.h

  Line 1 in 37724f7

  #ifndef DataFormats_SiPixelMappingSoA_alpaka_SiPixelClustersDevice_h

  • Change with:

   #ifndef CondFormats_SiPixelObjects_alpaka_SiPixelMappingDevice_h

  or

   #ifndef CondFormats_SiPixelObjects_interface_alpaka_SiPixelMappingDevice_h

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @fwyzard
  

  cmssw/CondFormats/SiPixelObjects/interface/alpaka/SiPixelMappingDevice.h

  Line 14 in 37724f7

  using SiPixelMappingHost = PortableHostCollection<SiPixelMappingLayout<>>;

  
  if this needed ?
  it's already defined in CondFormats/SiPixelObjects/interface/SiPixelMappingHost.h, outside of the ALPAKA_ACCELERATOR_NAMESPACE namespace.

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @fwyzard
  

  cmssw/DataFormats/BeamSpotAlpaka/interface/alpaka/BeamSpotAlpaka.h

  Line 29 in 37724f7

  cms::alpakatools::device_buffer<Device, BeamSpotPOD> const& ptr() const { return *data_d_; }

  • Change with:
    This is not safe: a const device_buffer can always be copied, and the copy used to modify the underlying data.
    Use a const_device_buffer instead:

       cms::alpakatools::const_device_buffer<Device, BeamSpotPOD> ptr() const { return *data_d_; }

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @fwyzard
  

  cmssw/DataFormats/BeamSpotAlpaka/BuildFile.xml

  Line 5 in 37724f7

  <flags ALPAKA_BACKENDS="serial"/>

  
  why only for the serial backend ?

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @fwyzard
  

  cmssw/DataFormats/BeamSpotAlpaka/BuildFile.xml

  Line 1 in 37724f7

  <use name="alpaka"/>

  
  I would prefer to avoid using Alpaka in the name of the types and of the packages.

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @fwyzard
  

  cmssw/DataFormats/BeamSpotAlpaka/interface/alpaka/BeamSpotAlpaka.h

  Line 11 in 37724f7

  class BeamSpotAlpaka {

  
  I would prefer to avoid using Alpaka in the name of the types and of the packages.

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @fwyzard
  

  cmssw/DataFormats/BeamSpotAlpaka/interface/alpaka/BeamSpotAlpaka.h

  Line 17 in 37724f7

  BeamSpotAlpaka(Queue& queue) : data_d_{cms::alpakatools::make_device_buffer<BeamSpotPOD>(queue)} {}

  • Change with:
    I think this should use a Queue const&:

       BeamSpotAlpaka(Queue const& queue) : data_d_{cms::alpakatools::make_device_buffer<BeamSpotPOD>(queue)} {}

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @fwyzard
  

  cmssw/DataFormats/BeamSpotAlpaka/src/alpaka/classes_serial_def.xml

  Line 1 in 37724f7

  <lcgdict>

  
  move this to DataFormats/BeamSpotAlpaka/src/classes_def.xml ?

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @fwyzard
  

  cmssw/DataFormats/BeamSpotAlpaka/src/alpaka/classes_serial.h

  Line 1 in 37724f7

  #ifndef DataFormats_BeamSpotAlpaka_alpaka_classes_serial_h

  
  move this to DataFormats/BeamSpotAlpaka/src/classes.h ?

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @fwyzard
  

  cmssw/DataFormats/Portable/interface/HostProductAlpaka.h

  Line 10 in 37724f7

  class HostProductAlpaka {

  
  This "heterogeneous unique pointer" should no longer be needed.
  Can you simply replace it simply with a cms::alpakatools::host_buffer<T> ?

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @fwyzard
  

  cmssw/DataFormats/TrackingRecHitSoA/interface/TrackingRecHitSoAHost.h

  Line 12 in 37724f7

  class TrackingRecHitAlpakaHost : public PortableHostCollection<TrackingRecHitAlpakaLayout<TrackerTraits>> {

  
  I would remove Alpaka from all type names

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @fwyzard
  

  cmssw/DataFormats/TrackingRecHitSoA/interface/TrackingRecHitsLayout.h

  Line 29 in 37724f7

  GENERATE_SOA_LAYOUT(TrackingRecHitAlpakaSoALayout,

  • Change with:
    As this is declared inside the struct TrackingRecHitAlpakaSoA, can it be called simply Layout ?

     GENERATE_SOA_LAYOUT(Layout,

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @fwyzard
  

  cmssw/DataFormats/TrackingRecHitSoA/interface/TrackingRecHitsLayout.h

  Line 46 in 37724f7

  //These above could be separated in a specific

  
  if the lifetime of the two "parts" is the same, I would keep them together

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @fwyzard
  

  cmssw/DataFormats/TrackingRecHitSoA/interface/TrackingRecHitsLayout.h

  Lines 57 to 59 in 37724f7

  	template <typename TrackerTraits>
  	using TrackingRecHitAlpakaLayout =
  	typename TrackingRecHitAlpakaSoA<TrackerTraits>::template TrackingRecHitAlpakaSoALayout<>;

  • Change with:
    I was going to suggest simpler names, like

   template <typename TrackerTraits>
   struct TrackingRecHitT {
     ...
     GENERATE_SOA_LAYOUT(Layout, ...);
   };
   
   template <typename TrackerTraits>
   using TrackingRecHitSoA = typename TrackingRecHitT<TrackerTraits>::template Layout<>;

  but I see that TrackingRecHitSoA is already used by the CUDA version :-/

  Any ideas for a naming scheme that does not include alpaka, and does not clash with the previous one ?

---

28 April 2023

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/CalibTracker/Records/interface/SiPixelMappingSoARecord.h

  Line 7 in 37724f7

  class SiPixelMappingSoARecord : public edm::eventsetup::EventSetupRecordImplementation<SiPixelMappingSoARecord> {};

  
  I see SiPixelRawToCluster consumes SiPixelMappingDevice in this record, but I don't see any ESProducer producing anything into SiPixelMappingSoARecord. I see SiPixelCablingSoAESProducer produces SiPixelMappingHost into CkfComponentsRecord that then gets automatically transferred into SiPixelMappingDevice. But which is the intention
  • Place the product in CkfComponentsRecord, in which case this Record and the consumer needs to be updated
  • Place the product in SiPixelMappingSoARecord, in which case this Record needs to be made to depend on some other Records, and the producer needs to be updated

  ?

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/CalibTracker/SiPixelESProducers/BuildFile.xml

  Line 13 in 37724f7

  <flags ALPAKA_BACKENDS="1"/>

  
  After moving the TYPELOOKUP_DATA_REG() registrations to their proper package, there should be no need to change anything in this BuildFile.xml.

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/CalibTracker/SiPixelESProducers/plugins/alpaka/SiPixelCablingSoAESProducer.cc

  Line 34 in 37724f7

  auto const& component = iConfig.getParameter<std::string>("ComponentName");

  
  I repeat my earlier comment that in principle this kind of ComponentName is redundant, as setting the data label could be achieved with the framework's built-in appendToDataLabel parameter that is available for all ESProducers.

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/CalibTracker/SiPixelESProducers/plugins/alpaka/SiPixelCablingSoAESProducer.cc

  Line 65 in 37724f7

  host()); //(pixelgpudetails::MAX_SIZE, *(cablingMap.product()), hasQuality, cms::alpakatools::host());

  
  Is the comment still necessary? If yes, consider moving to preceding line for better formatting.

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/CalibTracker/SiPixelESProducers/plugins/alpaka/SiPixelCablingSoAESProducer.cc

  Lines 70 to 71 in 37724f7

  	unsigned int startFed = *(fedIds.begin());
  	unsigned int endFed = *(fedIds.end() - 1);

  • Change with:
    How about

         unsigned int startFed = fedIds.front();
         unsigned int endFed = fedIds.back();

  ?

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/CalibTracker/SiPixelESProducers/plugins/alpaka/SiPixelGainCalibrationForHLTSoAESProducer.cc

  Line 44 in 37724f7

  //descriptions.add("siPixelGainCalibrationForHLTSoA", desc);

  
  Can this commented line be removed?

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/CalibTracker/SiPixelESProducers/plugins/alpaka/SiPixelGainCalibrationForHLTSoAESProducer.cc

  Lines 50 to 54 in 37724f7

  	auto gainsRecord = iRecord.getHandle(gainsToken_);
  	auto geomRecord = iRecord.getHandle(geometryToken_);
  	auto geom = *(geomRecord);
  	auto gains = *(gainsRecord);

  • Change with:
    Could be just

       auto const& gains= iRecord.get(gainsToken_);
       auto const& geom = iRecord.get(geometryToken_);

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/CalibTracker/SiPixelESProducers/plugins/alpaka/SiPixelGainCalibrationForHLTSoAESProducer.cc

  Lines 69 to 72 in 37724f7

  	//LogDebug("SiPixelGainCalibrationForHLTSoA")
  	std::cout << "caching calibs for " << n_detectors << " pixel detectors of size " << gains.data().size() << '\n'
  	<< "sizes " << sizeof(char) << ' ' << sizeof(uint8_t) << ' '
  	<< sizeof(siPixelGainsSoA::DecodingStructure);

  
  Switch back to LogDebug (or silence in some other way)

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/CalibTracker/SiPixelESProducers/plugins/alpaka/SiPixelGainCalibrationForHLTSoAESProducer.cc

  Line 74 in 37724f7

  // &(product->view().v_pedestals()) = (siPixelGainsSoA::DecodingStructure*)gains.data().data();

  
  Can this be removed?

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/CondFormats/SiPixelObjects/interface/alpaka/SiPixelGainCalibrationForHLTUtilities.h

  Line 1 in 37724f7
  • Suggestion:

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/CondFormats/SiPixelObjects/interface/alpaka/SiPixelMappingUtilities.h

  Line 1 in 37724f7
  • Suggestion:

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/DataFormats/SiPixelClusterSoA/interface/alpaka/SiPixelClustersDevice.h

  Line 16 in 37724f7

  class SiPixelClustersDevice : public PortableCollection<SiPixelClustersLayout<>> {

  
  Just to remind, this approach doesn't work for host backends, as for them the ALPAKA_ACCELERATOR_NAMESPACE::SiPixelClustersDevice must be the same type as SiPixelClustersHost.

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/DataFormats/BeamSpotAlpaka/src/alpaka/classes_cuda_def.xml

  Line 4 in 37724f7

  <class name="cms::alpakatools::Product<alpaka_cuda_async::BeamSpotAlpaka>" persistent="false"/>

  • Change with:

     <class name="edm::DeviceProduct<alpaka_cuda_async::BeamSpotAlpaka>" persistent="false"/>

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/DataFormats/BeamSpotAlpaka/src/alpaka/classes_rocm_def.xml

  Line 4 in 37724f7

  <class name="cms::alpakatools::Product<alpaka_rocm_async::BeamSpotAlpaka>" persistent="false"/>

  • Change with:

     <class name="edm::DeviceProduct<alpaka_rocm_async::BeamSpotAlpaka>" persistent="false"/>

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/DataFormats/SiPixelClusterSoA/src/alpaka/classes_cuda_def.xml

  Line 3 in 37724f7

  <class name="cms::alpakatools::Product<alpaka_cuda_async::Queue,alpaka_cuda_async::SiPixelClustersDevice>" persistent="false"/>

  • Change with:

     <class name="edm::DeviceProduct<alpaka_cuda_async::SiPixelClustersDevice>" persistent="false"/>

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/DataFormats/SiPixelClusterSoA/src/alpaka/classes_rocm_def.xml

  Line 3 in 37724f7

  <class name="cms::alpakatools::Product<alpaka_rocm_async::Queue,alpaka_rocm_async::SiPixelClustersDevice>" persistent="false"/>

  • Change with:

     <class name="edm::DeviceProduct<alpaka_rocm_async::SiPixelClustersDevice>" persistent="false"/>

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/DataFormats/SiPixelClusterSoA/src/alpaka/classes_serial_def.xml

  Line 1 in 37724f7

  <lcgdict>

  
  Move contents to DataFormats/SiPixelClusterSoA/src/classes_def.xml

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/DataFormats/SiPixelClusterSoA/src/alpaka/classes_serial.h

  Line 1 in 37724f7

  #ifndef DataFormats_SiPixelClusterSoA_alpaka_classes_serial_h

  
  Move contents to DataFormats/SiPixelClusterSoA/src/classes.h

---

05 May 2023

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/DataFormats/SiPixelDigiSoA/interface/SiPixelDigiErrorsHost.h

  Lines 20 to 21 in 37724f7

  	data_h = std::move(data);
  	error_h = cms::alpakatools::make_host_buffer<cms::alpakatools::SimpleVector<SiPixelErrorCompact>>();

  
  These could be done in the initializer list.

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/DataFormats/SiPixelDigiSoA/interface/SiPixelDigiErrorsHost.h

  Line 22 in 37724f7

  (*error_h).data()->set_data((*data_h).data());

  • Change with:
    How about

       error_h->data()->set_data(data_h->data());

  ? (also in general)

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/DataFormats/SiPixelDigiSoA/interface/SiPixelDigisHost.h

  Line 7 in 37724f7

  #include "SiPixelDigisLayout.h"

  • Change with:

   #include "DataFormats/SiPixelDigiSoA/interface/SiPixelDigisLayout.h"

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/DataFormats/SiPixelDigiSoA/interface/SiPixelDigisHost.h

  Lines 18 to 21 in 37724f7

  	~SiPixelDigisHost() = default;
  	SiPixelDigisHost(SiPixelDigisHost &&) = default;
  	SiPixelDigisHost &operator=(SiPixelDigisHost &&) = default;

  
  I'd suggest to either leave these to be declared by the compiler, or add the move operations explicitly.

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/DataFormats/SiPixelDigiSoA/interface/alpaka/SiPixelDigiErrorsDevice.h

  Lines 67 to 70 in 37724f7

  	SiPixelDigiErrorsHost dstData(error_vector.data()->capacity(), srcData.formatterErrors(), error_data_h);
  	// alpaka::memcpy(queue, dstData.buffer(), srcData.buffer());
  	return dstData;

  • Change with:
    How about

         return SiPixelDigiErrorsHost(error_vector.data()->capacity(), srcData.formatterErrors(), std::move(error_data_h));

  ?

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/DataFormats/SiPixelDigiSoA/interface/alpaka/SiPixelDigiErrorsDevice.h

  Line 59 in 37724f7

  auto error_vector = srcData.error_vector();

  • Change with:

         auto const& error_vector = srcData.error_vector();

  to avoid copy

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/DataFormats/SiPixelDigiSoA/interface/alpaka/SiPixelDigiErrorsDevice.h

  Line 61 in 37724f7

  auto error_data = srcData.error_data();

  • Change with:

         auto const& error_data = srcData.error_data();

  to avoid copy

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/DataFormats/SiPixelDigiSoA/interface/alpaka/SiPixelDigisDevice.h

  Lines 20 to 23 in 37724f7

  	~SiPixelDigisDevice() = default;
  	SiPixelDigisDevice(SiPixelDigisDevice &&) = default;
  	SiPixelDigisDevice &operator=(SiPixelDigisDevice &&) = default;

  
  I'd suggest to either rely on compiler declaring these, or add move oprations as well

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/DataFormats/SiPixelDigiSoA/src/alpaka/classes_cuda_def.xml

  Lines 3 to 5 in 37724f7

  	<class name="edm::Wrapper<alpaka_cuda_async::SiPixelDigisDevice>" persistent="false"/>
  	<class name="cms::alpakatools::Product<alpaka_cuda_async::Queue,alpaka_cuda_async::SiPixelDigisDevice>" persistent="false"/>
  	<class name="edm::Wrapper<cms::alpakatools::Product<alpaka_cuda_async::Queue,alpaka_cuda_async::SiPixelDigisDevice>>" persistent="false"/>

  • Change with:
    Should be

     <class name="edm::DeviceProduct<alpaka_cuda_async::SiPixelDigisDevice>" persistent="false"/>
     <class name="edm::Wrapper<edm::DeviceProduct<alpaka_cuda_async::SiPixelDigisDevice>>" persistent="false"/>

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/DataFormats/SiPixelDigiSoA/src/alpaka/classes_rocm_def.xml

  Lines 3 to 5 in 37724f7

  	<class name="edm::Wrapper<alpaka_rocm_async::SiPixelDigisDevice>" persistent="false"/>
  	<class name="cms::alpakatools::Product<alpaka_rocm_async::Queue,alpaka_rocm_async::SiPixelDigisDevice>" persistent="false"/>
  	<class name="edm::Wrapper<cms::alpakatools::Product<alpaka_rocm_async::Queue,alpaka_rocm_async::SiPixelDigisDevice>>" persistent="false"/>

  • Change with:
    Should be

     <class name="edm::DeviceProduct<alpaka_rocm_async::SiPixelDigisDevice>" persistent="false"/>
     <class name="edm::Wrapper<edm::DeviceProduct<alpaka_rocm_async::SiPixelDigisDevice>>" persistent="false"/>

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/DataFormats/SiPixelDigiSoA/src/alpaka/classes_serial.h

  Line 1 in 37724f7

  #ifndef DataFormats_SiPixelDigiSoA_alpaka_alasses_serial_h

  
  This file should be removed

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/DataFormats/SiPixelDigiSoA/src/alpaka/classes_serial_def.xml

  Line 1 in 37724f7

  <lcgdict>

  
  This file should be removed

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/DataFormats/SiPixelDigiSoA/interface/alpaka/SiPixelDigisDevice.h

  Lines 12 to 14 in 37724f7

  	namespace ALPAKA_ACCELERATOR_NAMESPACE {
  	class SiPixelDigisDevice : public PortableCollection<SiPixelDigisLayout<>> {

  
  Needs the type alias for alpaka_serial_sync::SiPixelDigisDevice = SiPixelDigisHost

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/DataFormats/SiPixelDigiSoA/interface/alpaka/SiPixelDigiErrorsDevice.h

  Lines 15 to 17 in 37724f7

  	namespace ALPAKA_ACCELERATOR_NAMESPACE {
  	class SiPixelDigiErrorsDevice {

  
  Needs type alias for alpaka_serial_sync::SiPixelDigiErrorsDevice = SiPixelDigiErrorsHost

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/DataFormats/SiPixelDigiSoA/src/classes_def.xml

  Line 17 in 37724f7

  <class name="SiPixelDigisHost"/>

  
  I believe this needs class version

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/DataFormats/SiPixelDigiSoA/src/classes_def.xml

  Line 19 in 37724f7

  <class name="SiPixelDigiErrorsHost"/>

  
  I beliveve this needs class version

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/DataFormats/SiPixelClusterSoA/src/classes_def.xml

  Line 17 in 37724f7

  <class name="SiPixelClustersHost"/>

  
  I believe this needs class version

---

12 May 2023

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/DataFormats/TrackingRecHitSoA/interface/alpaka/TrackingRecHitSoADevice.h

  Line 85 in 37724f7

  using TrackingRecHitAlpakaDevicePhase2 = TrackingRecHitAlpakaDevice<pixelTopology::Phase2>;

  
  Need type aliases

   alpaka_serial_sync::TrackingRecHitAlpakaDevicePhase1 = TrackingRecHitAlpakaHostPhase1;
   alpaka_serial_sync::TrackingRecHitAlpakaDevicePhase2 = TrackingRecHitAlpakaHostPhase2;
  

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/DataFormats/TrackingRecHitSoA/interface/alpaka/TrackingRecHitSoADevice.h

  Line 101 in 37724f7

  struct CopyToHost<ALPAKA_ACCELERATOR_NAMESPACE::TrackingRecHitAlpakaDevicePhase2> {

  
  You could avoid copy-paste by templating this along

   template <typename TrackerTraits>
   struct CopyToHost<ALPAKA_ACCELERATOR_NAMESPACE::TrackingRecHitAlpakaDevice<TrackerTraits> {
   ...

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/DataFormats/TrackingRecHitSoA/src/alpaka/classes_cuda_def.xml

  Lines 3 to 5 in 37724f7

  	<class name="edm::Wrapper<alpaka_cuda_async::TrackingRecHitAlpakaDevicePhase1>" persistent="false"/>
  	<class name="cms::alpakatools::Product<alpaka_cuda_async::Queue,alpaka_cuda_async::TrackingRecHitAlpakaDevicePhase1>" persistent="false"/>
  	<class name="edm::Wrapper<cms::alpakatools::Product<alpaka_cuda_async::Queue,alpaka_cuda_async::TrackingRecHitAlpakaDevicePhase1>>" persistent="false"/>

  • Change with:

     <class name="edm::DeviceProduct<alpaka_cuda_async::TrackingRecHitAlpakaDevicePhase1>" persistent="false"/>
     <class name="edm::Wrapper<edm::DeviceProduct<alpaka_cuda_async::TrackingRecHitAlpakaDevicePhase1>>" persistent="false"/>

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/DataFormats/TrackingRecHitSoA/src/alpaka/classes_cuda_def.xml

  Lines 3 to 5 in 37724f7

  	<class name="edm::Wrapper<alpaka_cuda_async::TrackingRecHitAlpakaDevicePhase1>" persistent="false"/>
  	<class name="cms::alpakatools::Product<alpaka_cuda_async::Queue,alpaka_cuda_async::TrackingRecHitAlpakaDevicePhase1>" persistent="false"/>
  	<class name="edm::Wrapper<cms::alpakatools::Product<alpaka_cuda_async::Queue,alpaka_cuda_async::TrackingRecHitAlpakaDevicePhase1>>" persistent="false"/>

  • Change with:

     <class name="edm::DeviceProduct<alpaka_cuda_async::TrackingRecHitAlpakaDevicePhase1>" persistent="false"/>
     <class name="edm::Wrapper<edm::DeviceProduct<alpaka_cuda_async::TrackingRecHitAlpakaDevicePhase1>>" persistent="false"/>

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/DataFormats/TrackingRecHitSoA/src/alpaka/classes_cuda_def.xml

  Lines 8 to 10 in 37724f7

  	<class name="edm::Wrapper<alpaka_cuda_async::TrackingRecHitAlpakaDevicePhase2>" persistent="false"/>
  	<class name="cms::alpakatools::Product<alpaka_cuda_async::Queue,alpaka_cuda_async::TrackingRecHitAlpakaDevicePhase2>" persistent="false"/>
  	<class name="edm::Wrapper<cms::alpakatools::Product<alpaka_cuda_async::Queue,alpaka_cuda_async::TrackingRecHitAlpakaDevicePhase2>>" persistent="false"/>

  • Change with:

     <class name="edm::DeviceProduct<alpaka_cuda_async::TrackingRecHitAlpakaDevicePhase2>" persistent="false"/>
     <class name="edm::Wrapper<edm::DeviceProduct<alpaka_cuda_async::TrackingRecHitAlpakaDevicePhase2>>" persistent="false"/>

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/DataFormats/TrackingRecHitSoA/src/alpaka/classes_rocm_def.xml

  Lines 3 to 5 in 37724f7

  	<class name="edm::Wrapper<alpaka_rocm_async::TrackingRecHitAlpakaDevicePhase1>" persistent="false"/>
  	<class name="cms::alpakatools::Product<alpaka_rocm_async::Queue,alpaka_rocm_async::TrackingRecHitAlpakaDevicePhase1>" persistent="false"/>
  	<class name="edm::Wrapper<cms::alpakatools::Product<alpaka_rocm_async::Queue,alpaka_rocm_async::TrackingRecHitAlpakaDevicePhase1>>" persistent="false"/>

  • Change with:

     <class name="edm::DeviceProduct<alpaka_rocm_async::TrackingRecHitAlpakaDevicePhase1>" persistent="false"/>
     <class name="edm::Wrapper<edm::DeviceProduct<alpaka_rocm_async::TrackingRecHitAlpakaDevicePhase1>>" persistent="false"/>

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/DataFormats/TrackingRecHitSoA/src/alpaka/classes_rocm_def.xml

  Lines 8 to 10 in 37724f7

  	<class name="edm::Wrapper<alpaka_rocm_async::TrackingRecHitAlpakaDevicePhase2>" persistent="false"/>
  	<class name="cms::alpakatools::Product<alpaka_rocm_async::Queue,alpaka_rocm_async::TrackingRecHitAlpakaDevicePhase2>" persistent="false"/>
  	<class name="edm::Wrapper<cms::alpakatools::Product<alpaka_rocm_async::Queue,alpaka_rocm_async::TrackingRecHitAlpakaDevicePhase2>>" persistent="false"/>

  • Change with:

     <class name="edm::DeviceProduct<alpaka_rocm_async::TrackingRecHitAlpakaDevicePhase2>" persistent="false"/>
     <class name="edm::Wrapper<edm::DeviceProduct<alpaka_rocm_async::TrackingRecHitAlpakaDevicePhase2>>" persistent="false"/>

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/DataFormats/TrackingRecHitSoA/src/alpaka/classes_serial.h

  Line 1 in 37724f7

  #ifndef DataFormats_TrackingRecHitSoA_alpaka_classes_serial_h

  
  Remove this file

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/DataFormats/TrackingRecHitSoA/src/alpaka/classes_serial_def.xml

  Line 1 in 37724f7

  <lcgdict>

  
  Remove this file

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/DataFormats/TrackingRecHitSoA/src/classes_def.xml

  Line 22 in 37724f7

  <class name="TrackingRecHitAlpakaHostPhase1"/>

  
  I believe this class would need a version

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/DataFormats/TrackingRecHitSoA/src/classes_def.xml

  Line 37 in 37724f7

  <class name="TrackingRecHitAlpakaHostPhase2"/>

  
  I believe this class would need a version

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/DataFormats/TrackingRecHitSoA/interface/TrackingRecHitSoAHost.h

  Line 43 in 37724f7

  ParamsOnDevice const* cpeParams_;

  
  Is cpeParams_ supposed to be persisted? (I don't actually see it set or used anywhere)

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/DataFormats/TrackingRecHitSoA/interface/TrackingRecHitSoAHost.h

  Line 46 in 37724f7

  PhiBinnerStorageType* phiBinnerStorage_;

  
  Is phiBinnerStorage_ supposed to be persisted?

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/DataFormats/TrackingRecHitSoA/interface/TrackingRecHitsLayout.h

  Line 9 in 37724f7

  #include "SiPixelHitStatus.h"

  • Change with:

   #include "DataFormats/TrackingRecHitSoA/interface/SiPixelHitStatus.h"

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/DataFormats/TrackingRecHitSoA/test/BuildFile.xml

  Line 5 in 37724f7

  <flags ALPAKA_BACKENDS="1"/>

  • Change with:

     <flags ALPAKA_BACKENDS="1"/>

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/DataFormats/TrackingRecHitSoA/src/classes_def.xml

  Line 10 in 37724f7

  <class name="PortableHostCollection<TrackingRecHitAlpakaLayout<pixelTopology::Phase1>>"/>

  • Change with:

     <class name="PortableHostCollection<TrackingRecHitAlpakaLayout<pixelTopology::Phase1>>"/>

---

16 May 2023

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/DataFormats/TrackingRecHitSoA/test/alpaka/Hits_test.cc

  Line 48 in 37724f7

  */

  
  This test is not very useful with the body being commented out. Does it not work, or is there some other complication?

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelClusterizer/BuildFile.xml

  Line 1 in 37724f7

  <use name="FWCore/Framework"/>

  
  This file does not seem to be needed (because there is no code in src or interface directories)

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelClusterizer/plugins/BuildFile.xml

  Line 27 in 37724f7

  <use name="HeterogeneousCore/AlpakaTest"/>

  • Suggestion: The `AlpakaTes package should not be needed

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/CalibPixel.h

  Line 26 in 37724f7

  ; // must be > MaxNumModules

  
  Spurious semicolon

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/CalibPixel.h

  Line 85 in 37724f7

  // });

  
  Does this commented out code mean that the implementation is still incomplete?

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/CalibPixel.h

  Line 22 in 37724f7

  namespace ALPAKA_ACCELERATOR_NAMESPACE {

  
  Because of calibPixel::operator() being templated over TAcc, I believe this code does not need to be enclosed in ALPAKA_ACCELERATOR_NAMESPACE.

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/PixelClustering.h

  Line 27 in 37724f7

  struct PixelStatus {

  
  It looks like this could be a namespace?

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/PixelClustering.h

  Line 141 in 37724f7

  acc); // packed words array used to store the PixelStatus of each pixel

  
  Would moving comment above the variable definition improve formatting?

  Also, auto& status should work equally well, and be less confusing.

---

17 May 2023

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/PixelClustering.h

  Lines 270 to 275 in 37724f7

  	#if defined(ALPAKA_ACC_GPU_CUDA_ASYNC_BACKEND) || defined(ALPAKA_ACC_GPU_HIP_ASYNC_BACKEND)
  	constexpr uint32_t threadDimension = 1;
  	#else
  	// NB: can be tuned.
  	constexpr uint32_t threadDimension = 256;
  	#endif

  • Change with:
    We should generally avoid customizing behavior per backend with macros in functions that are (effectively) templated over the backend. How about something along

         if constexpr (std::is_same_v<alpaka::DevCpu, alpaka::Dev<TAcc>>) {
           // NB: can be tuned.
           constexpr uint32_t threadDimension = 256;
         } else {
           constexpr uint32_t threadDimension = 1;
         }

  ?

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/PixelClustering.h

  Lines 448 to 449 in 37724f7

  	#endif
  	#ifdef GPU_DEBUG

  • Suggestion: Maybe just merge the two #ifdef blocks?

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/SiPixelPhase2DigiToCluster.cc

  Lines 18 to 24 in 37724f7

  	#include "FWCore/Framework/interface/ConsumesCollector.h"
  	#include "FWCore/Framework/interface/ESHandle.h"
  	#include "FWCore/Framework/interface/ESTransientHandle.h"
  	#include "FWCore/Framework/interface/ESWatcher.h"
  	#include "FWCore/Framework/interface/Event.h"
  	#include "FWCore/Framework/interface/EventSetup.h"
  	#include "FWCore/Framework/interface/stream/EDProducer.h"

  • Suggestion: These includes could be removed

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/SiPixelPhase2DigiToCluster.cc

  Line 29 in 37724f7

  #include "FWCore/ServiceRegistry/interface/Service.h"

  • Suggestion:
    not needed

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/SiPixelPhase2DigiToCluster.cc

  Line 31 in 37724f7

  #include "HeterogeneousCore/AlpakaCore/interface/ScopedContext.h"

  • Suggestion:
    doesn't exist anymore

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/SiPixelPhase2DigiToCluster.cc

  Lines 34 to 36 in 37724f7

  	#include "FWCore/ParameterSet/interface/ConfigurationDescriptions.h"
  	#include "FWCore/ParameterSet/interface/ParameterSet.h"
  	#include "FWCore/ParameterSet/interface/ParameterSetDescription.h"

  • Suggestion:
    already included above

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/SiPixelPhase2DigiToCluster.cc

  Lines 43 to 44 in 37724f7

  	#include "HeterogeneousCore/AlpakaTest/interface/AlpakaESTestRecords.h"
  	#include "HeterogeneousCore/AlpakaTest/interface/alpaka/AlpakaESTestData.h"

  • Suggestion:
    should not be needed in production code

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/SiPixelPhase2DigiToCluster.cc

  Line 63 in 37724f7

  // cms::alpakatools::ContextState<Queue> ctxState_;

  • Suggestion:

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/SiPixelPhase2DigiToCluster.cc

  Lines 104 to 105 in 37724f7

  	// cms::cuda::ScopedContextAcquire ctx{iEvent.streamID(), std::move(waitingTaskHolder), ctxState_};

  • Suggestion:

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/SiPixelRawToCluster.cc

  Lines 27 to 29 in 37724f7

  	#include "FWCore/Framework/interface/ConsumesCollector.h"
  	#include "FWCore/Framework/interface/ESHandle.h"
  	#include "FWCore/Framework/interface/ESTransientHandle.h"

  • Suggestion:

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/SiPixelRawToCluster.cc

  Lines 31 to 33 in 37724f7

  	#include "FWCore/Framework/interface/Event.h"
  	#include "FWCore/Framework/interface/EventSetup.h"
  	#include "FWCore/Framework/interface/stream/EDProducer.h"

  • Suggestion:

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/SiPixelRawToCluster.cc

  Line 38 in 37724f7

  #include "FWCore/ServiceRegistry/interface/Service.h"

  • Suggestion:

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/SiPixelRawToCluster.cc

  Line 40 in 37724f7

  #include "HeterogeneousCore/AlpakaCore/interface/ScopedContext.h"

  • Suggestion:

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/SiPixelRawToCluster.cc

  Lines 43 to 45 in 37724f7

  	#include "FWCore/ParameterSet/interface/ConfigurationDescriptions.h"
  	#include "FWCore/ParameterSet/interface/ParameterSet.h"
  	#include "FWCore/ParameterSet/interface/ParameterSetDescription.h"

  • Suggestion:

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/SiPixelRawToCluster.cc

  Lines 52 to 53 in 37724f7

  	#include "HeterogeneousCore/AlpakaTest/interface/AlpakaESTestRecords.h"
  	#include "HeterogeneousCore/AlpakaTest/interface/alpaka/AlpakaESTestData.h"

  • Suggestion:

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/SiPixelRawToCluster.cc

  Line 71 in 37724f7

  // cms::alpakatools::ContextState<Queue> ctxState_;

  • Suggestion:

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/SiPixelRawToCluster.cc

  Lines 112 to 113 in 37724f7

  	digiErrorPutToken_ =
  	produces(); //<cms::alpakatools::Product<alpaka_cuda_async::Queue,alpaka_cuda_async::SiPixelDigisDevice>>(); //reg.produces<cms::alpakatools::Product<Queue, SiPixelDigiErrorsAlpaka>>();

  • Change with:

         digiErrorPutToken_ = produces();

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/SiPixelRawToCluster.cc

  Line 99 in 37724f7

  : rawGetToken_(consumes<FEDRawDataCollection>(iConfig.getParameter<edm::InputTag>("InputLabel"))),

  
  Having a mixture of type-deducing consumes() and explicit-type consumes() looks a bit strange. How about using the type-deducing consumes for all of the input data products?

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/SiPixelRawToCluster.cc

  Lines 148 to 149 in 37724f7

  	// cms::alpakatools::ScopedContextAcquire<Queue> ctx{iEvent.streamID(), std::move(waitingTaskHolder), ctxState_};

  • Suggestion:

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/SiPixelRawToCluster.cc

  Lines 156 to 158 in 37724f7

  	// get the GPU product already here so that the async transfer can begin
  	// const auto* Map = hMap.getGPUProductAsync(iEvent.queue());
  	// const unsigned char* ModulesToUnpack = hMap.getModToUnpAllAsync(iEvent.queue());

  • Suggestion:

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/SiPixelRawToCluster.cc

  Line 170 in 37724f7

  cablingMap_ = iSetup.getHandle(cablingMapToken_).product();

  • Change with:

         cablingMap_ = &iSetup.getData(cablingMapToken_);

  would be better

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/SiPixelRawToCluster.cc

  Line 168 in 37724f7

  if (recordWatcher_.check(iSetup) or regions_) {

  
  Is the or regions_ part in the condition necessary? As far as I can tell, the fedIds_ is not changed by other code, so updating it at IOV boundaries should be sufficient?

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/SiPixelRawToCluster.cc

  Line 184 in 37724f7

  *(regions_->modulesToUnpack()), (cablingMap_->cablingTree()).get(), cablingMap_->fedIds(), iEvent.queue());

  • Change with:
    How about using the cached values?

             *(regions_->modulesToUnpack()), cabling_.get(), fedIds_, iEvent.queue());

  ?

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/SiPixelRawToCluster.cc

  Lines 294 to 297 in 37724f7

  	SiPixelDigisDevice digis_d = SiPixelDigisDevice(nDigis_, iEvent.queue());
  	SiPixelClustersDevice clusters_d = SiPixelClustersDevice(pixelTopology::Phase1::numberOfModules, iEvent.queue());
  	iEvent.emplace(digiPutToken_, std::move(digis_d));
  	iEvent.emplace(clusterPutToken_, std::move(clusters_d));

  • Change with:
    Could be just

         iEvent.emplace(digiPutToken_, nDigis_, iEvent.queue());
         iEvent.emplace(clusterPutToken_, pixelTopology::Phase1::numberOfModules, iEvent.queue());

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/SiPixelRawToCluster.cc

  Line 299 in 37724f7

  iEvent.emplace(digiErrorPutToken_, SiPixelDigiErrorsDevice());

  • Change with:
    Could be just

           iEvent.emplace(digiErrorPutToken_);

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/SiPixelRawToCluster.cc

  Lines 304 to 305 in 37724f7

  	// auto tmp = Algo_.getResults();
  	iEvent.emplace(digiPutToken_, Algo_.getDigis()); //std::move(tmp.first));

  • Change with:

       iEvent.emplace(digiPutToken_, Algo_.getDigis());

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/SiPixelRawToClusterKernel.dev.cc

  Line 362 in 37724f7

  // const SiPixelFedCablingMapGPU *cablingMap,

  • Suggestion:

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/SiPixelRawToClusterKernel.dev.cc

  Lines 369 to 374 in 37724f7

  	// uint16_t *xx,
  	// uint16_t *yy,
  	// uint16_t *adc,
  	// uint32_t *pdigi,
  	// uint32_t *rawIdArr,
  	// uint16_t *moduleId,

  • Suggestion:

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/SiPixelRawToClusterKernel.dev.cc

  Line 494 in 37724f7

  struct fillHitsModuleStart {

  
  It looks a bit strange to have one kernel defined inside ALPAKA_ACCELERATOR_NAMESPACE and others outside of it.

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/SiPixelRawToClusterKernel.dev.cc

  Lines 497 to 499 in 37724f7

  	// uint32_t const *__restrict__ cluStart,
  	SiPixelClustersLayoutSoAView clus_view
  	// uint32_t *__restrict__ moduleStart

  • Change with:

                                     SiPixelClustersLayoutSoAView clus_view

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/SiPixelRawToClusterKernel.dev.cc

  Lines 584 to 587 in 37724f7

  	// const SiPixelFedCablingMapGPU *cablingMap,
  	const SiPixelMappingLayoutSoAConstView &cablingMap,
  	const unsigned char *modToUnp,
  	// const SiPixelGainForHLTonGPU *gains,

  • Change with:

                                                         const SiPixelMappingLayoutSoAConstView &cablingMap,
                                                         const unsigned char *modToUnp,

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/SiPixelRawToClusterKernel.dev.cc

  Lines 612 to 617 in 37724f7

  	#if defined(ALPAKA_ACC_GPU_CUDA_ASYNC_BACKEND) || defined(ALPAKA_ACC_GPU_HIP_ASYNC_BACKEND)
  	const int threadsPerBlockOrElementsPerThread = 512;
  	#else
  	// NB: MPORTANT: This could be tuned to benefit from innermost loop.
  	const int threadsPerBlockOrElementsPerThread = 32;
  	#endif

  • Change with:
    Or how about

           const int threadsPerBlockOrElementsPerThread = []() {
             if constexpr (std::is_same_v<Device, alpaka::DevHost>) {
               // NB: MPORTANT: This could be tuned to benefit from innermost loop.
               return 32;
             } else {
               return 512;
             }
           }();

  (although we'd probably benefit from some general mechanism for specifying device/backend-specific launch parameters)

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/SiPixelRawToClusterKernel.dev.cc

  Lines 634 to 636 in 37724f7

  	alpaka::enqueue(queue,
  	alpaka::createTaskKernel<Acc1D>(workDiv,
  	RawToDigi_kernel(),

  • Change with:
    How about

           alpaka::exec<Acc1D>(queue,
                               RawToDigi_kernel(),

  ?

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/SiPixelRawToClusterKernel.dev.cc

  Lines 643 to 648 in 37724f7

  	// digis_d->xx(),
  	// digis_d->yy(),
  	// digis_d->adc(),
  	// digis_d->pdigi(),
  	// digis_d->rawIdArr(),
  	// digis_d->moduleInd(),

  • Change with:

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/SiPixelRawToClusterKernel.dev.cc

  Line 660 in 37724f7

  digiErrors_d.copyErrorToHostAsync(stream);

  
  Just to note this is not done yet

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/SiPixelRawToClusterKernel.dev.cc

  Line 671 in 37724f7

  #if defined(ALPAKA_ACC_GPU_CUDA_ASYNC_BACKEND) || defined(ALPAKA_ACC_GPU_HIP_ASYNC_BACKEND)

  
  if constexpr instead?

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/SiPixelRawToClusterKernel.dev.cc

  Lines 681 to 687 in 37724f7

  	// alpaka::enqueue(queue, alpaka::createTaskKernel<Acc1D>(workDiv, calibDigis(),
  	// isRun2,
  	// digis_d->view(),
  	// clusters_d->view(),
  	// gains,
  	// wordCounter));

  • Suggestion:

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/SiPixelRawToClusterKernel.dev.cc

  Line 694 in 37724f7

  alpaka::enqueue(

  
  alpaka::exec() instead?

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/SiPixelRawToClusterKernel.dev.cc

  Line 713 in 37724f7

  alpaka::enqueue(queue,

  
  alpaka::exec() instead?

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/SiPixelRawToClusterKernel.dev.cc

  Lines 719 to 724 in 37724f7

  	// template <typename TAcc>
  	// ALPAKA_FN_ACC void operator()(
  	// const TAcc& acc,
  	// SiPixelDigisLayoutSoAView digi_view,
  	// SiPixelClustersLayoutSoAView clus_view,
  	// const unsigned int numElements)

  • Change with:

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/SiPixelRawToClusterKernel.dev.cc

  Line 731 in 37724f7

  alpaka::enqueue(queue,

  
  alpaka::exec()?

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/SiPixelRawToClusterKernel.dev.cc

  Line 746 in 37724f7

  alpaka::enqueue(

  
  alpaka::exec()?

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/SiPixelRawToClusterKernel.dev.cc

  Lines 778 to 780 in 37724f7

  	// digis_d = SiPixelDigisDevice(numDigis, queue);
  	// alpaka::memcpy(queue, digis_d->view(), digis_h_view);

  • Suggestion:

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/SiPixelRawToClusterKernel.dev.cc

  Line 784 in 37724f7

  // nModules_Clusters_h = cms::cuda::make_host_unique<uint32_t[]>(2, stream);

  • Suggestion:

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/SiPixelRawToClusterKernel.dev.cc

  Lines 791 to 793 in 37724f7

  	// alpaka::enqueue(queue,
  	// alpaka::createTaskKernel<Acc1D>(
  	// workDiv, calibPixel::calibDigisPhase2{}, digis_view, clusters_d->view(), numDigis));

  • Suggestion:

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/SiPixelRawToClusterKernel.dev.cc

  Line 800 in 37724f7

  alpaka::enqueue(queue,

  
  alpaka::exec()?

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/SiPixelRawToClusterKernel.dev.cc

  Line 815 in 37724f7

  alpaka::enqueue(

  
  alpaka::exec()?

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/SiPixelRawToClusterKernel.dev.cc

  Line 824 in 37724f7

  alpaka::enqueue(queue,

  
  alpaka::exec()?

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/SiPixelRawToClusterKernel.dev.cc

  Line 839 in 37724f7

  alpaka::enqueue(

  
  alpaka::exec()?

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/SiPixelRawToClusterKernel.dev.cc

  Line 843 in 37724f7

  // clusters_d->clusModuleStart()));

  • Suggestion:

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/SiPixelRawToClusterKernel.h

  Lines 23 to 24 in 37724f7

  	// struct SiPixelFedCablingMapGPU;
  	// class SiPixelGainForHLTonGPU;

  • Suggestion:

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/SiPixelRawToClusterKernel.h

  Lines 168 to 172 in 37724f7

  	// std::pair<SiPixelDigisDevice, SiPixelClustersDevice> getResults() {
  	// digis_d->setNModulesDigis(nModules_Clusters_h[0], nDigis);
  	// clusters_d->setNClusters(nModules_Clusters_h[1], nModules_Clusters_h[2]);
  	// return std::make_pair(std::move(*digis_d), std::move(*clusters_d));
  	// }

  • Suggestion:

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/SiPixelRawToClusterKernel.h

  Line 174 in 37724f7

  SiPixelDigisDevice&& getDigis() {

  • Change with:
    The proper return type is just

         SiPixelDigisDevice getDigis() {

  (same for SiPixelClustersDevice and SiPixelDigiErrorsDevice below)

---

18 May 2023

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelRecHits/BuildFile.xml

  Lines 12 to 13 in 37724f7

  	<use name="CUDADataFormats/TrackingRecHit"/>
  	<use name="CUDADataFormats/BeamSpot"/>

  
  Why is dependence on CUDA data formats added in this PR?

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelRecHits/interface/PixelCPEFastParams.h

  Lines 153 to 154 in 37724f7

  	template <>
  	struct CopyToDevice<pixelCPEforDevice::PixelCPEFastParams<alpaka_common::DevHost, pixelTopology::Phase1>> {

  • Change with:
    The two full specializations could be provided with a partial specialization along

     template <typename TrackerTraits>
     struct CopyToDevice<pixelCPEforDevice::PixelCPEFastParams<alpaka_common::DevHost, TrackerTraits>> {

---

19 May 2023

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelRecHits/plugins/BuildFile.xml

  Line 12 in 37724f7

  <use name="CUDADataFormats/TrackingRecHit"/>

  
  Could all the CUDA dependencies be moved under <library file="*.cc ${cuda_src}" name="RecoLocalTrackerSiPixelRecHitsPlugins">?
  (in principle all the dependencies not used by the Alpaka code would be good to move)

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitFromSoAAlpaka.cc

  Lines 25 to 26 in 37724f7

  	#include "HeterogeneousCore/AlpakaInterface/interface/config.h"
  	#include "HeterogeneousCore/AlpakaInterface/interface/memory.h"

  
  Are these two includes really needed?

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitFromSoAAlpaka.cc

  Lines 95 to 96 in 37724f7

  	// Could not make a view on hitsModuleStart
  	//alpaka::memcpy(iEvent.queue(), hmsp, hitsModuleStartView);

  • Suggestion: Not needed?

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelRecHits/plugins/alpaka/PixelCPEFastParamsESProducerAlpaka.cc

  Line 30 in 37724f7

  auto const& myname = iConfig.getParameter<std::string>("ComponentName");

  
  This line does not do anything. Actually the whole ESProducer is a no-op because the constructor does not register the production function with setWhatProduced(). What does it try to do?

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelRecHits/plugins/alpaka/PixelRecHitGPUKernel.dev.cc

  Line 6 in 37724f7

  //#include <alpaka/alpaka.hpp>

  • Change with:

   #include <alpaka/alpaka.hpp>

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelRecHits/plugins/alpaka/PixelRecHitGPUKernel.dev.cc

  Lines 80 to 81 in 37724f7

  	// getHits<TrackerTraits><<<blocks, threadsPerBlock, 0, stream>>>(
  	// cpeParams, bs_d.data(), digis_d.view(), digis_d.nDigis(), clusters_d.const_view(), hits_d.view());

  • Suggestion:

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelRecHits/plugins/alpaka/PixelRecHitGPUKernel.dev.cc

  Lines 89 to 90 in 37724f7

  	// setHitsLayerStart<TrackerTraits>
  	// <<<1, 32, 0, stream>>>(clusters_d->clusModuleStart(), cpeParams, hits_d.view().hitsLayerStart().data());

  • Suggestion:

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelRecHits/plugins/alpaka/SiPixelRecHitAlpaka.cc

  Line 5 in 37724f7

  #include "DataFormats/Common/interface/Handle.h"

  • Suggestion:

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelRecHits/plugins/alpaka/SiPixelRecHitAlpaka.cc

  Line 78 in 37724f7

  throw cms::Exception("Configuration") << "SiPixelRecHitAlpaka can only use a CPE of type PixelCPEFast";

  
  How does the nullness of fcpe.data() tell about the type of the CPE?

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelRecHits/interface/PixelCPEFastParams.h

  Line 133 in 37724f7

  class PixelCPEFastParams {

  
  This class seems to be only a wrapper for a buffer, and the actual data product of interest is the object held by the buffer. Would it be useful if the EventSetup system would support a case where the ESProducer can produce the host-side data product as alpaka host buffer (instead of std::unique_ptr etc), the system would implicitly copy it to am alpaka device buffer, and the consuming module would consume the consume the product directly (i.e. in this case the consumer would use a token of type device::ESGetToken<pixelCPEforDevice::ParamsOnDeviceT<TrackerTraits>, TkPixelCPERecord>)?

  (I'm thinking longer term evolution beyond this CUDA->Alpaka migration, i.e. I would not change this part of this PR)

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelRecHits/interface/PixelCPEFastAlpaka.h

  Line 16 in 37724f7

  class PixelCPEFastAlpaka final : public PixelCPEGenericBase {

  
  It looks like this is an EventSetup data product, but I didn't see any ESProducer for it.

• Pixel Alpaka Migration: Local Reconstruction [III] cms-sw/cmssw#41285 (comment) by @makortel
  

  cmssw/RecoLocalTracker/SiPixelRecHits/src/PixelCPEFastAlpaka.cc

  Lines 345 to 354 in 37724f7

  	// template <typename TrackerTraits>
  	// PixelCPEFastAlpaka<TrackerTraits>::GPUData::~GPUData() {
  	// if (paramsOnGPU_d != nullptr) {
  	// cudaFree((void*)paramsOnGPU_h.m_commonParams);
  	// cudaFree((void*)paramsOnGPU_h.m_detParams);
  	// cudaFree((void*)paramsOnGPU_h.m_averageGeometry);
  	// cudaFree((void*)paramsOnGPU_h.m_layerGeometry);
  	// cudaFree(paramsOnGPU_d);
  	// }
  	// }

  • Suggestion: