start working on validation

RecoLocalCalo/HGCalRecProducers/plugins/HGCalRecHitKernelImpl.cu

@@ -171,7 +171,7 @@ void fill_positions_from_detids(const hgcal_conditions::HeterogeneousHEFConditio
       xpos += (1.5f * (cV - ncells) + 1.f) * R1;
       ypos += (cU - 0.5f * cV - n2) * r1_x2;
 
-      conds->posmap.x[i] = xpos; //* side; multiply by -1 if one wants to obtain the position from the opposite endcap
+      conds->posmap.x[i] = xpos; //* side; multiply by -1 if one wants to obtain the position from the opposite endcap. CAREFUL WITH LATER DETECTOR ALIGNMENT!!!
       conds->posmap.y[i] = ypos;
 
       //printf( "%d - %lf - %lf\n", cV - ncells, 1.5f*(static_cast<float>(cV) - static_cast<float>(ncells)), 1.5f*(static_cast<float>(cV - ncells)) );
@@ -188,7 +188,7 @@ void print_positions_from_detids(const hgcal_conditions::HeterogeneousHEFConditi
     {
       HeterogeneousHGCSiliconDetId did(conds->posmap.detid[i]);
       const int32_t layer = did.layer();
-      float posz = conds->posmap.z_per_layer[ layer-1 ]; //multiply by -1 if one wants to obtain the position from the opposite endcap
+      float posz = conds->posmap.z_per_layer[ layer-1 ];
       printf("PosX: %lf\t PosY: %lf\t Posz: %lf\n", conds->posmap.x[i], conds->posmap.y[i], posz);
     } 
 }

RecoLocalCalo/HGCalRecProducers/plugins/HeterogeneousHGCalHEFRecHitProducer.cc

@@ -135,6 +135,7 @@ void HeterogeneousHGCalHEFRecHitProducer::set_conditions_(const edm::EventSetup&
 
   int upper_estimate_wafer_number  = 2 * ddd_->lastLayer(true) * (ddd_->waferMax() - ddd_->waferMin());
   int upper_estimate_cell_number = upper_estimate_wafer_number * 24 * 24; 
+  posmap_->z_per_layer.resize( (ddd_->lastLayer(true) - ddd_->firstLayer()) * 2 );
   posmap_->numberCellsHexagon.reserve(upper_estimate_wafer_number);
   posmap_->detid.reserve(upper_estimate_cell_number);
   //set positons-related variables
@@ -148,7 +149,8 @@ void HeterogeneousHGCalHEFRecHitProducer::set_conditions_(const edm::EventSetup&
 
   //store detids following a geometry ordering
   for(int ilayer=1; ilayer<=posmap_->lastLayer; ++ilayer) {
-    posmap_->z_per_layer.push_back( static_cast<float>( ddd_->waferZ(ilayer, true) ) ); //originally a double
+    posmap_->z_per_layer[ilayer-1+(ddd_->lastLayer(true) - ddd_->firstLayer())] = static_cast<float>( ddd_->waferZ(ilayer, true) ); //originally a double
+    posmap_->z_per_layer[ilayer-1] = static_cast<float>( ddd_->waferZ(-ilayer, true) ); //originally a double
 
     for(int iwaferU=posmap_->waferMin; iwaferU<posmap_->waferMax; ++iwaferU) {
       for(int iwaferV=posmap_->waferMin; iwaferV<posmap_->waferMax; ++iwaferV) {

RecoLocalCalo/HGCalRecProducers/plugins/KernelManagerHGCalRecHit.cu

@@ -52,17 +52,18 @@ void KernelManagerHGCalRecHit::run_kernels(const KernelConstantData<HGChefUncali
 {
   transfer_soas_to_device_();
 
-  fill_positions_from_detids<<<::nblocks_,::nthreads_>>>(d_conds);
-  after_();
+  //fill_positions_from_detids<<<::nblocks_,::nthreads_>>>(d_conds);
+  //after_();
 
-  print_positions_from_detids<<<::nblocks_,::nthreads_>>>(d_conds);
+  //print_positions_from_detids<<<::nblocks_,::nthreads_>>>(d_conds);
   /*
   hef_step1<<<::nblocks_,::nthreads_>>>( *(data_->d_2), *(data_->d_1_), d_kcdata->data, data_->nhits_);
   after_();
   reuse_device_pointers_();
   */
 
-  //hef_to_rechit<<<::nblocks_,::nthreads_>>>( *(data_->d_out_), *(data_->d_1_), kcdata->data_, d_conds, data_->nhits_ );
+  after_();
+  hef_to_rechit<<<::nblocks_,::nthreads_>>>( *(data_->d_out_), *(data_->d_1_), kcdata->data_, d_conds, data_->nhits_ );
   after_();
 
   transfer_soa_to_host_and_synchronize_();

UserCode/RecHitsValidator/BuildFile.xml

@@ -0,0 +1,7 @@
+<use name="root"/>
+<use name="RecoLocalCalo/HGCalRecProducers"/>
+<export>
+  <lib   name="1"/>
+</export>
+<Flags CXXFLAGS="-O0"/>
+<Flags CXXFLAGS="-g"/>

UserCode/RecHitsValidator/plugins/BuildFile.xml

@@ -0,0 +1,15 @@
+<library name="RecHitsValidator" file="HeterogeneousHGCalRecHitsValidator.cc">
+  <use name="cuda"/>
+  <use name="HeterogeneousCore/CUDACore"/>
+  <use name="HeterogeneousCore/CUDAUtilities"/>
+  <use name="DataFormats/HcalDetId"/>
+  <use name="FWCore/Framework"/>
+  <use name="FWCore/MessageLogger"/>
+  <use name="FWCore/Utilities"/>
+  <use name="RecoLocalCalo/HGCalRecAlgos"/>
+  <use name="RecoLocalCalo/HGCalRecProducers"/>
+  <use name="Geometry/HGCalGeometry"/>
+  <use name="Geometry/HGCalCommonData"/>
+  <use name="UserCode/RecHitsValidator"/>
+  <flags EDM_PLUGIN="1"/>
+</library>

UserCode/RecHitsValidator/plugins/HeterogeneousHGCalRecHitsValidator.cc

@@ -0,0 +1,84 @@
+#include "UserCode/RecHitsValidator/plugins/HeterogeneousHGCalRecHitsValidator.h"
+
+#include "DetectorDescription/OfflineDBLoader/interface/GeometryInfoDump.h"
+#include "Geometry/Records/interface/IdealGeometryRecord.h"
+#include "Geometry/HGCalGeometry/interface/HGCalGeometry.h"
+#include "SimG4CMS/Calo/interface/CaloHitID.h"
+
+#include "DetectorDescription/Core/interface/DDFilter.h"
+#include "DetectorDescription/Core/interface/DDFilteredView.h"
+#include "DetectorDescription/Core/interface/DDSolid.h"
+
+#include "DataFormats/GeometryVector/interface/Basic3DVector.h"
+
+#include "CLHEP/Units/GlobalSystemOfUnits.h"
+
+#include "FWCore/ServiceRegistry/interface/Service.h"
+#include "FWCore/Utilities/interface/Exception.h"
+#include "CommonTools/UtilAlgos/interface/TFileService.h"
+
+HeterogeneousHGCalRecHitsValidator::HeterogeneousHGCalRecHitsValidator( const edm::ParameterSet &ps ) : 
+  tokens_( {{ {{consumes<HGCRecHitCollection>(ps.getParameter<edm::InputTag>("cpuRecHitsCEEToken")),
+		consumes<HGCRecHitCollection>(ps.getParameter<edm::InputTag>("gpuRecHitsCEEToken"))}},
+	      {{consumes<HGCRecHitCollection>(ps.getParameter<edm::InputTag>("cpuRecHitsCEEToken")),
+		consumes<HGCRecHitCollection>(ps.getParameter<edm::InputTag>("gpuRecHitsCEEToken"))}},
+	      {{consumes<HGCRecHitCollection>(ps.getParameter<edm::InputTag>("cpuRecHitsCEEToken")),
+		consumes<HGCRecHitCollection>(ps.getParameter<edm::InputTag>("gpuRecHitsCEEToken"))}} }} ),
+  treename_("tree")
+{
+  edm::Service<TFileService> fs;
+  tree_ = fs->make<TTree>(treename_.c_str(), treename_.c_str());
+  tree_->Branch( "cpuDetid", "std::vector<unsigned int>", &cpuValidRecHits.detid );
+  tree_->Branch( "gpuDetid", "std::vector<unsigned int>", &gpuValidRecHits.detid );
+  //zhist.push_back(fs->make<TH1F>(("z"+treename_).c_str(), ("z"+treename_).c_str(), 2000, -10, 650));
+}
+
+HeterogeneousHGCalRecHitsValidator::~HeterogeneousHGCalRecHitsValidator()
+{
+}
+
+void HeterogeneousHGCalRecHitsValidator::endJob()
+{
+}
+
+void HeterogeneousHGCalRecHitsValidator::set_geometry_(const edm::EventSetup& setup, const unsigned int& detidx)
+{
+  edm::ESHandle<HGCalGeometry> handle;
+  setup.get<IdealGeometryRecord>().get(handles_str_[detidx], handle);
+}
+
+void HeterogeneousHGCalRecHitsValidator::analyze(const edm::Event &event, const edm::EventSetup &setup)
+{
+  recHitTools_.getEventSetup(setup);
+
+  //future subdetector loop
+  for(size_t idet=0; idet<1; ++idet) {
+    set_geometry_(setup, 1/*idet*/);
+
+    //get hits produced with the CPU
+    event.getByToken(tokens_[idet][0], handles_[idet][0]);
+    const auto &cpuhits = *handles_[idet][0];
+
+    //get hits produced with the GPU
+    event.getByToken(tokens_[idet][1], handles_[idet][1]);
+    const auto &gpuhits = *handles_[idet][1];
+
+    size_t nhits = cpuhits.size();
+    assert( nhits == gpuhits.size() );
+    for(unsigned int i=0; i<nhits; i++) {
+      const HGCRecHit &cpuHit = cpuhits[i];
+      const HGCRecHit &gpuHit = gpuhits[i];
+
+      const HGCalDetId cpuDetid = cpuHit.detid();
+      const HGCalDetId gpuDetid = gpuHit.detid();
+      cpuValidRecHits.detid.push_back( cpuDetid );
+      gpuValidRecHits.detid.push_back( gpuDetid );
+
+      std::cout << cpuDetid << ", " << gpuDetid << std::endl;
+    }
+  }    
+  tree_->Fill();
+}
+
+//define this as a plug-in
+DEFINE_FWK_MODULE(HeterogeneousHGCalRecHitsValidator);

UserCode/RecHitsValidator/plugins/Validator.h → UserCode/RecHitsValidator/plugins/HeterogeneousHGCalRecHitsValidator.h

@@ -33,6 +33,11 @@
 #include <iostream>
 #include <string>
 
+struct ValidRecHits {
+  std::vector<float> son;
+  std::vector<unsigned int> detid;
+};
+
 class HeterogeneousHGCalRecHitsValidator : public edm::EDAnalyzer 
 {
  public:
@@ -42,22 +47,20 @@ class HeterogeneousHGCalRecHitsValidator : public edm::EDAnalyzer
   void endJob();
 
  private:
-  edm::EDGetTokenT<edm::HepMCProduct> mc_;
-  std::vector< edm::EDGetTokenT<HGCRecHitCollection> > recHitsTokens_;
-  edm::EDGetTokenT<reco::GenParticleCollection> genParticles_;
-
+  static const unsigned int nsubdetectors = 3; //ce-e, ce-h-fine, ce-h-coarse
+  static const unsigned int ncomputingdevices = 2; //cpu, gpu
+  //cpu amd gpu tokens and handles for the 3 subdetectors, cpu and gpu
+  std::array< std::array< edm::EDGetTokenT<HGChefRecHitCollection>, ncomputingdevices>, nsubdetectors> tokens_; 
+  std::array< std::array< edm::Handle<HGChefRecHitCollection>, ncomputingdevices>, nsubdetectors> handles_;
+  std::array< std::string, nsubdetectors> handles_str_ = {{"HGCalEESensitive", "HGCalHESiliconSensitive", "HGCalHEScintillatorSensitive"}};
   hgcal::RecHitTools recHitTools_;
 
-  std::vector< TH1F* > zhist;
-  std::vector< TH1F* > xhist;
-  std::vector< TH1F* > yhist; 
-  std::vector< TH1F* > zsidehist;
-  std::vector< TH1F* > ehist;
-  std::vector< TH1F* > layerhist;
-  std::vector< TH1F* > offsetlayerhist;
-
   TTree* tree_;
   std::string treename_;
+  ValidRecHits cpuValidRecHits, gpuValidRecHits;
+  //std::vector< TH1F* > zhist;
+
+  void set_geometry_(const edm::EventSetup&, const unsigned int&);
 };
 
 

UserCode/RecHitsValidator/python/HeterogeneousHGCalRecHitsValidator_cfg.py

@@ -0,0 +1,70 @@
+import os, sys, glob
+import FWCore.ParameterSet.Config as cms
+from Configuration.StandardSequences.Eras import eras
+from Configuration.ProcessModifiers.gpu_cff import gpu
+from Configuration.AlCa.GlobalTag import GlobalTag
+process.GlobalTag = GlobalTag(process.GlobalTag, 'auto:phase2_realistic', '')
+
+#package loading
+process = cms.Process("gpuValidation", gpu) 
+process.load('FWCore.MessageService.MessageLogger_cfi')
+process.load('Configuration.StandardSequences.Services_cff')
+process.load('SimGeneral.HepPDTESSource.pythiapdt_cfi')
+process.load('Configuration.StandardSequences.MagneticField_cff')
+#process.load('Configuration.EventContent.EventContent_cff')
+process.load('Configuration.StandardSequences.FrontierConditions_GlobalTag_cff')
+process.load('Configuration.Geometry.GeometryExtended2026D46Reco_cff')
+process.load('HeterogeneousCore.CUDAServices.CUDAService_cfi')
+process.load('RecoLocalCalo.HGCalRecProducers.HGCalRecHit_cfi')
+process.load('SimCalorimetry.HGCalSimProducers.hgcalDigitizer_cfi')
+
+from RecoLocalCalo.HGCalRecProducers.HGCalRecHit_cfi import dEdX_weights_v10
+process.maxEvents = cms.untracked.PSet( input = cms.untracked.int32(-1) )
+
+#TFileService
+fileName = 'validation.root'
+process.TFileService = cms.Service("TFileService", 
+                                   fileName = cms.string(fileName),
+                                   closeFileFast = cms.untracked.bool(True)
+                               )
+
+
+fNames = ['file:/afs/cern.ch/user/b/bfontana/CMSSW_11_1_0_pre6/src/20495.0_CloseByPGun_CE_E_Front_200um+CE_E_Front_200um_2026D41_GenSimHLBeamSpotFull+DigiFullTrigger_2026D41+RecoFullGlobal_2026D41+HARVESTFullGlobal_2026D41/step3.root']
+keep = 'keep *'
+drop = 'drop CSCDetIdCSCALCTPreTriggerDigiMuonDigiCollection_simCscTriggerPrimitiveDigis__HLT'
+process.source = cms.Source("PoolSource",
+                            fileNames = cms.untracked.vstring(fNames),
+                            inputCommands = cms.untracked.vstring([keep, drop]),
+                            duplicateCheckMode = cms.untracked.string("noDuplicateCheck"))
+
+process.options = cms.untracked.PSet(
+    wantSummary = cms.untracked.bool( False )) #add option for edmStreams
+process.HeterogeneousHGCalHEFRecHits = cms.EDProducer('HeterogeneousHGCalHEFRecHitProducer',
+                                                      HGCHEFUncalibRecHitsTok = cms.InputTag('HGCalUncalibRecHit', 'HGCHEFUncalibRecHits'),
+                                                      HGCHEF_keV2DIGI  = HGCalRecHit.__dict__['HGCHEF_keV2DIGI'],
+                                                      minValSiPar     = HGCalRecHit.__dict__['minValSiPar'],
+                                                      maxValSiPar     = HGCalRecHit.__dict__['maxValSiPar'],
+                                                      constSiPar      = HGCalRecHit.__dict__['constSiPar'],
+                                                      noiseSiPar      = HGCalRecHit.__dict__['noiseSiPar'],
+                                                      HGCHEF_fCPerMIP = HGCalRecHit.__dict__['HGCHEF_fCPerMIP'],
+                                                      HGCHEF_isSiFE   = HGCalRecHit.__dict__['HGCHEF_isSiFE'],
+                                                      HGCHEF_noise_fC = HGCalRecHit.__dict__['HGCHEF_noise_fC'],
+                                                      HGCHEF_cce      = HGCalRecHit.__dict__['HGCHEF_cce'],
+                                                      rangeMatch      = HGCalRecHit.__dict__['rangeMatch'],
+                                                      rangeMask       = HGCalRecHit.__dict__['rangeMask'],
+                                                      rcorr           = HGCalRecHit.__dict__['thicknessCorrection'],
+                                                      weights         = HGCalRecHit.__dict__['layerWeights'] )
+
+
+process.valid = cms.EDAnalyzer( "HeterogeneousHGCalRecHitsValidator",
+                                cpuRecHitsHSiToken = cms.InputTag('HeterogeneousHGCalHEFRecHits','HGCHEFRecHits'),
+                                gpuRecHitsHSiToken = cms.InputTag('HeterogeneousHGCalHEFRecHits','HGCHEFRecHits') )
+
+
+process.task = cms.Task( process.HeterogeneousHGCalHEFRecHitProducer * process.valid )
+process.path = cms.Path( process.task )
+
+
+process.out = cms.OutputModule( "PoolOutputModule", 
+                                fileName = cms.untracked.string('out.root') )
+process.outpath = cms.EndPath(process.out)