stream syncrhonization finalized

CondFormats/HGCalObjects/src/HeterogeneousHGCalHEFCellPositionsConditions.cc

@@ -94,11 +94,11 @@ std::vector<size_t> HeterogeneousHGCalHEFCellPositionsConditions::calculate_memo
   for(unsigned int i=0; i<npointers; ++i)
     {
       if(cpos::types[i] == cpos::HeterogeneousHGCalPositionsType::Float)
-	sizes[i] = select_pointer_u_(cpuPos, 1).size(); //each position array (x, y and z) will have the same size as the detid array
+	sizes[i] = select_pointer_u_(cpuPos, 2).size(); //each position array (x, y and z) will have the same size as the detid array
       else if(cpos::types[i] == cpos::HeterogeneousHGCalPositionsType::Int32_t)
-	sizes[i] = select_pointer_i_(cpuPos, 0).size();
+	sizes[i] = select_pointer_i_(cpuPos, 1).size();
       else if(cpos::types[i] == cpos::HeterogeneousHGCalPositionsType::Uint32_t)
-	sizes[i] = select_pointer_u_(cpuPos, 1).size();
+	sizes[i] = select_pointer_u_(cpuPos, 2).size();
     }
 
   std::vector<size_t> sizes_units(npointers);

RecoLocalCalo/HGCalRecProducers/plugins/HGCalRecHitKernelImpl.cu

@@ -121,7 +121,7 @@ void ee_to_rechit(HGCRecHitSoA dst_soa, HGCUncalibratedRecHitSoA src_soa, const
 }
 
 __global__
-void hef_to_rechit(HGCRecHitSoA dst_soa, HGCUncalibratedRecHitSoA src_soa, const HGChefUncalibratedRecHitConstantData cdata, const hgcal_conditions::HeterogeneousHEFConditionsESProduct* conds, int length)
+void hef_to_rechit(HGCRecHitSoA dst_soa, HGCUncalibratedRecHitSoA src_soa, const HGChefUncalibratedRecHitConstantData cdata, int length)
 {
   unsigned int tid = blockDim.x * blockIdx.x + threadIdx.x;
 

RecoLocalCalo/HGCalRecProducers/plugins/HGCalRecHitKernelImpl.cuh

@@ -24,7 +24,7 @@ __global__
 void ee_to_rechit(HGCRecHitSoA dst_soa, HGCUncalibratedRecHitSoA src_soa, HGCeeUncalibratedRecHitConstantData cdata, int length);
 
 __global__
-void hef_to_rechit(HGCRecHitSoA dst_soa, HGCUncalibratedRecHitSoA src_soa, HGChefUncalibratedRecHitConstantData cdata, const hgcal_conditions::HeterogeneousHEFConditionsESProduct*, int length);
+void hef_to_rechit(HGCRecHitSoA dst_soa, HGCUncalibratedRecHitSoA src_soa, HGChefUncalibratedRecHitConstantData cdata, int length);
 
 __global__
 void heb_to_rechit(HGCRecHitSoA dst_soa, HGCUncalibratedRecHitSoA src_soa, HGChebUncalibratedRecHitConstantData cdata, int length);

RecoLocalCalo/HGCalRecProducers/plugins/HeterogeneousHGCalEERecHitProducer.cc

@@ -34,15 +34,15 @@ std::string HeterogeneousHGCalEERecHitProducer::assert_error_message_(std::strin
 
 void HeterogeneousHGCalEERecHitProducer::assert_sizes_constants_(const HGCConstantVectorData& vd)
 {
-  if( vdata_.fCPerMIP_.size() != maxsizes_constants::ee_fCPerMIP )
+  if( vdata_.fCPerMIP_.size() > maxsizes_constants::ee_fCPerMIP )
     cms::cuda::LogError("WrongSize") << this->assert_error_message_("fCPerMIP", maxsizes_constants::ee_fCPerMIP, vdata_.fCPerMIP_.size());
-  else if( vdata_.cce_.size() != maxsizes_constants::ee_cce )
+  else if( vdata_.cce_.size() > maxsizes_constants::ee_cce )
     cms::cuda::LogError("WrongSize") << this->assert_error_message_("cce", maxsizes_constants::ee_cce, vdata_.cce_.size());
-  else if( vdata_.noise_fC_.size() != maxsizes_constants::ee_noise_fC )
+  else if( vdata_.noise_fC_.size() > maxsizes_constants::ee_noise_fC )
     cms::cuda::LogError("WrongSize") << this->assert_error_message_("noise_fC", maxsizes_constants::ee_noise_fC, vdata_.noise_fC_.size());
-  else if( vdata_.rcorr_.size() != maxsizes_constants::ee_rcorr ) 
+  else if( vdata_.rcorr_.size() > maxsizes_constants::ee_rcorr ) 
     cms::cuda::LogError("WrongSize") << this->assert_error_message_("rcorr", maxsizes_constants::ee_rcorr, vdata_.rcorr_.size());
-  else if( vdata_.weights_.size() != maxsizes_constants::ee_weights ) 
+  else if( vdata_.weights_.size() > maxsizes_constants::ee_weights ) 
     cms::cuda::LogError("WrongSize") << this->assert_error_message_("weights", maxsizes_constants::ee_weights, vdata_.weights_.size());
 }
 
@@ -59,12 +59,13 @@ void HeterogeneousHGCalEERecHitProducer::acquire(edm::Event const& event, edm::E
     {
       allocate_memory_(ctx.stream());
       kcdata_ = new KernelConstantData<HGCeeUncalibratedRecHitConstantData>(cdata_, vdata_);
+      kmdata_ = new KernelModifiableData<HGCUncalibratedRecHitSoA, HGCRecHitSoA>(nhits, stride_, uncalibSoA_, d_uncalibSoA_, d_intermediateSoA_, d_calibSoA_, calibSoA_);
+
       convert_constant_data_(kcdata_);
       convert_collection_data_to_soa_(hits_ee, uncalibSoA_, nhits);
-      kmdata_ = new KernelModifiableData<HGCUncalibratedRecHitSoA, HGCRecHitSoA>(nhits, stride_, uncalibSoA_, d_uncalibSoA_, d_intermediateSoA_, d_calibSoA_, calibSoA_);
 
       KernelManagerHGCalRecHit kernel_manager(kmdata_);
-      kernel_manager.run_kernels(kcdata_);
+      kernel_manager.run_kernels(kcdata_, ctx.stream());
 
       rechits_ = std::make_unique<HGCRecHitCollection>();
       convert_soa_data_to_collection_(*rechits_, calibSoA_, nhits);

RecoLocalCalo/HGCalRecProducers/plugins/HeterogeneousHGCalHEBRecHitProducer.cc

@@ -27,7 +27,7 @@ std::string HeterogeneousHGCalHEBRecHitProducer::assert_error_message_(std::stri
 
 void HeterogeneousHGCalHEBRecHitProducer::assert_sizes_constants_(const HGCConstantVectorData& vd)
 {
-  if( vdata_.weights_.size() != maxsizes_constants::heb_weights )
+  if( vdata_.weights_.size() > maxsizes_constants::heb_weights )
     cms::cuda::LogError("MaxSizeExceeded") << this->assert_error_message_("weights", vdata_.fCPerMIP_.size());
 }
 
@@ -61,12 +61,13 @@ void HeterogeneousHGCalHEBRecHitProducer::acquire(edm::Event const& event, edm::
     {
       allocate_memory_(ctx.stream());
       kcdata_ = new KernelConstantData<HGChebUncalibratedRecHitConstantData>(cdata_, vdata_);  
+      kmdata_ = new KernelModifiableData<HGCUncalibratedRecHitSoA, HGCRecHitSoA>(nhits, stride_, uncalibSoA_, d_uncalibSoA_, d_intermediateSoA_, d_calibSoA_, calibSoA_);
+
       convert_constant_data_(kcdata_);
       convert_collection_data_to_soa_(hits_heb, uncalibSoA_, nhits);
-      kmdata_ = new KernelModifiableData<HGCUncalibratedRecHitSoA, HGCRecHitSoA>(nhits, stride_, uncalibSoA_, d_uncalibSoA_, d_intermediateSoA_, d_calibSoA_, calibSoA_);
 
       KernelManagerHGCalRecHit kernel_manager(kmdata_);
-      kernel_manager.run_kernels(kcdata_);
+      kernel_manager.run_kernels(kcdata_, ctx.stream());
 
       convert_soa_data_to_collection_(*rechits_, calibSoA_, nhits);
       deallocate_memory_();

RecoLocalCalo/HGCalRecProducers/plugins/HeterogeneousHGCalHEFRecHitProducer.cc

@@ -45,15 +45,15 @@ std::string HeterogeneousHGCalHEFRecHitProducer::assert_error_message_(std::stri
 
 void HeterogeneousHGCalHEFRecHitProducer::assert_sizes_constants_(const HGCConstantVectorData& vd)
 {
-  if( vdata_.fCPerMIP_.size() != maxsizes_constants::hef_fCPerMIP )
+  if( vdata_.fCPerMIP_.size() > maxsizes_constants::hef_fCPerMIP )
     cms::cuda::LogError("WrongSize") << this->assert_error_message_("fCPerMIP", maxsizes_constants::hef_fCPerMIP, vdata_.fCPerMIP_.size());
-  else if( vdata_.cce_.size() != maxsizes_constants::hef_cce )
+  else if( vdata_.cce_.size() > maxsizes_constants::hef_cce )
     cms::cuda::LogError("WrongSize") << this->assert_error_message_("cce", maxsizes_constants::hef_cce, vdata_.cce_.size());
-  else if( vdata_.noise_fC_.size() != maxsizes_constants::hef_noise_fC )
+  else if( vdata_.noise_fC_.size() > maxsizes_constants::hef_noise_fC )
     cms::cuda::LogError("WrongSize") << this->assert_error_message_("noise_fC", maxsizes_constants::hef_noise_fC, vdata_.noise_fC_.size());
-  else if( vdata_.rcorr_.size() != maxsizes_constants::hef_rcorr )
+  else if( vdata_.rcorr_.size() > maxsizes_constants::hef_rcorr )
     cms::cuda::LogError("WrongSize") << this->assert_error_message_("rcorr", maxsizes_constants::hef_rcorr, vdata_.rcorr_.size());
-  else if( vdata_.weights_.size() != maxsizes_constants::hef_weights ) 
+  else if( vdata_.weights_.size() > maxsizes_constants::hef_weights ) 
     cms::cuda::LogError("WrongSize") << this->assert_error_message_("weights", maxsizes_constants::hef_weights, vdata_.weights_.size());
 }
 
@@ -100,12 +100,13 @@ void HeterogeneousHGCalHEFRecHitProducer::acquire(edm::Event const& event, edm::
     {
       allocate_memory_(ctx.stream());
       kcdata_ = new KernelConstantData<HGChefUncalibratedRecHitConstantData>(cdata_, vdata_);
+      kmdata_ = new KernelModifiableData<HGCUncalibratedRecHitSoA, HGCRecHitSoA>(nhits, stride_, uncalibSoA_, d_uncalibSoA_, d_intermediateSoA_, d_calibSoA_, calibSoA_);
+
       convert_constant_data_(kcdata_);
       convert_collection_data_to_soa_(hits_hef, uncalibSoA_, nhits);
-      kmdata_ = new KernelModifiableData<HGCUncalibratedRecHitSoA, HGCRecHitSoA>(nhits, stride_, uncalibSoA_, d_uncalibSoA_, d_intermediateSoA_, d_calibSoA_, calibSoA_);
 
       KernelManagerHGCalRecHit kernel_manager(kmdata_);
-      kernel_manager.run_kernels(kcdata_, nullptr/*d_conds*/);
+      kernel_manager.run_kernels(kcdata_, ctx.stream());
 
       convert_soa_data_to_collection_(*rechits_, calibSoA_, nhits);
       deallocate_memory_(); //cannot be called in destructor, since pointers are created in a conditional expression

RecoLocalCalo/HGCalRecProducers/plugins/KernelManagerHGCalRecHit.cu

@@ -21,15 +21,15 @@ KernelManagerHGCalRecHit::~KernelManagerHGCalRecHit()
 {
 }
 
-void KernelManagerHGCalRecHit::transfer_soas_to_device_()
+void KernelManagerHGCalRecHit::transfer_soas_to_device_(const cudaStream_t& stream)
 {
-  cudaCheck( cudaMemcpyAsync((data_->d_1_)->amplitude_, (data_->h_in_)->amplitude_, nbytes_device_, cudaMemcpyHostToDevice) );
+  cudaCheck( cudaMemcpyAsync((data_->d_1_)->amplitude_, (data_->h_in_)->amplitude_, nbytes_device_, cudaMemcpyHostToDevice, stream) );
   cudaCheck( cudaGetLastError() );
 }
 
-void KernelManagerHGCalRecHit::transfer_soa_to_host_and_synchronize_()
+void KernelManagerHGCalRecHit::transfer_soa_to_host_and_synchronize_(const cudaStream_t& stream)
 {
-  cudaCheck( cudaMemcpyAsync((data_->h_out_)->energy_, (data_->d_out_)->energy_, nbytes_host_, cudaMemcpyDeviceToHost) );
+  cudaCheck( cudaMemcpyAsync((data_->h_out_)->energy_, (data_->d_out_)->energy_, nbytes_host_, cudaMemcpyDeviceToHost, stream) );
   cudaCheck( cudaGetLastError() );
 }
 
@@ -38,26 +38,28 @@ void KernelManagerHGCalRecHit::reuse_device_pointers_()
   std::swap(data_->d_1_, data_->d_2_); 
 }
 
-void KernelManagerHGCalRecHit::run_kernels(const KernelConstantData<HGCeeUncalibratedRecHitConstantData> *kcdata)
+void KernelManagerHGCalRecHit::run_kernels(const KernelConstantData<HGCeeUncalibratedRecHitConstantData> *kcdata, const cudaStream_t& stream)
 {
-  transfer_soas_to_device_();
+  transfer_soas_to_device_( stream );
+  cudaCheck( cudaGetLastError() );
+
   /*
-  ee_step1<<<::nblocks_, ::nthreads_>>>( *(data_->d_2_), *(data_->d_1_), kcdata->data_, data_->nhits_ );
+  ee_step1<<<::nblocks_, ::nthreads_, 0, stream>>>( *(data_->d_2_), *(data_->d_1_), kcdata->data_, data_->nhits_ );
   after_();
   reuse_device_pointers_();
   */
 
-  ee_to_rechit<<<::nblocks_, ::nthreads_>>>( *(data_->d_out_), *(data_->d_1_), kcdata->data_, data_->nhits_ );
+  ee_to_rechit<<<::nblocks_, ::nthreads_, 0, stream>>>( *(data_->d_out_), *(data_->d_1_), kcdata->data_, data_->nhits_ );
   cudaCheck( cudaGetLastError() );
 
-  transfer_soa_to_host_and_synchronize_();
+  transfer_soa_to_host_and_synchronize_( stream );
   cudaCheck( cudaGetLastError() );
-  cudaCheck( cudaDeviceSynchronize() );
+  cudaCheck( cudaStreamSynchronize(stream) );
 }
 
-void KernelManagerHGCalRecHit::run_kernels(const KernelConstantData<HGChefUncalibratedRecHitConstantData> *kcdata, const hgcal_conditions::HeterogeneousHEFConditionsESProduct* d_conds)
+void KernelManagerHGCalRecHit::run_kernels(const KernelConstantData<HGChefUncalibratedRecHitConstantData> *kcdata, const cudaStream_t& stream)
 {
-  transfer_soas_to_device_();
+  transfer_soas_to_device_( stream );
   cudaCheck( cudaGetLastError() );
 
   /*
@@ -66,17 +68,17 @@ void KernelManagerHGCalRecHit::run_kernels(const KernelConstantData<HGChefUncali
   reuse_device_pointers_();
   */
 
-  hef_to_rechit<<<::nblocks_,::nthreads_>>>( *(data_->d_out_), *(data_->d_1_), kcdata->data_, d_conds, data_->nhits_ );
+  hef_to_rechit<<<::nblocks_, ::nthreads_, 0, stream>>>( *(data_->d_out_), *(data_->d_1_), kcdata->data_, data_->nhits_ );
   cudaCheck( cudaGetLastError() );
 
-  transfer_soa_to_host_and_synchronize_();
+  transfer_soa_to_host_and_synchronize_( stream );
   cudaCheck( cudaGetLastError() );
-  cudaCheck( cudaDeviceSynchronize() );
+  cudaCheck( cudaStreamSynchronize(stream) );
 }
 
-void KernelManagerHGCalRecHit::run_kernels(const KernelConstantData<HGChebUncalibratedRecHitConstantData> *kcdata)
+void KernelManagerHGCalRecHit::run_kernels(const KernelConstantData<HGChebUncalibratedRecHitConstantData> *kcdata, const cudaStream_t& stream)
 {
-  transfer_soas_to_device_();
+  transfer_soas_to_device_( stream );
   cudaCheck( cudaGetLastError() );
 
   /*
@@ -85,12 +87,12 @@ void KernelManagerHGCalRecHit::run_kernels(const KernelConstantData<HGChebUncali
   reuse_device_pointers_();
   */
 
-  heb_to_rechit<<<::nblocks_, ::nthreads_>>>( *(data_->d_out_), *(data_->d_1_), kcdata->data_, data_->nhits_ );
+  heb_to_rechit<<<::nblocks_, ::nthreads_, 0, stream>>>( *(data_->d_out_), *(data_->d_1_), kcdata->data_, data_->nhits_ );
   cudaCheck( cudaGetLastError() );
 
-  transfer_soa_to_host_and_synchronize_();
+  transfer_soa_to_host_and_synchronize_( stream );
   cudaCheck( cudaGetLastError() );
-  cudaCheck( cudaDeviceSynchronize() );
+  cudaCheck( cudaStreamSynchronize(stream) );
 }
 
 void KernelManagerHGCalRecHit::fill_positions(const hgcal_conditions::HeterogeneousHEFCellPositionsConditionsESProduct* d_conds)

RecoLocalCalo/HGCalRecProducers/plugins/KernelManagerHGCalRecHit.h

@@ -55,16 +55,15 @@ class KernelManagerHGCalRecHit {
   KernelManagerHGCalRecHit();
   KernelManagerHGCalRecHit(KernelModifiableData<HGCUncalibratedRecHitSoA, HGCRecHitSoA>*);
   ~KernelManagerHGCalRecHit();
-  void run_kernels(const KernelConstantData<HGCeeUncalibratedRecHitConstantData>*);
-  void run_kernels(const KernelConstantData<HGChefUncalibratedRecHitConstantData>*, const hgcal_conditions::HeterogeneousHEFConditionsESProduct*);
-  void run_kernels(const KernelConstantData<HGChebUncalibratedRecHitConstantData>*);
+  void run_kernels(const KernelConstantData<HGCeeUncalibratedRecHitConstantData>*, const cudaStream_t&);
+  void run_kernels(const KernelConstantData<HGChefUncalibratedRecHitConstantData>*, const cudaStream_t&);
+  void run_kernels(const KernelConstantData<HGChebUncalibratedRecHitConstantData>*, const cudaStream_t&);
   void fill_positions(const hgcal_conditions::HeterogeneousHEFCellPositionsConditionsESProduct*);
   HGCRecHitSoA* get_output();
 
  private:
-  void after_();
-  void transfer_soas_to_device_();
-  void transfer_soa_to_host_and_synchronize_();
+  void transfer_soas_to_device_(const cudaStream_t&);
+  void transfer_soa_to_host_and_synchronize_(const cudaStream_t&);
   void reuse_device_pointers_();
 
   int nbytes_host_;

UserCode/CodeGPU/python/HeterogeneousHGCalRecHit_cfg.py

@@ -1,4 +1,5 @@
 import FWCore.ParameterSet.Config as cms
+import os, glob
 
 enableGPU = True
 from Configuration.ProcessModifiers.gpu_cff import gpu
@@ -24,9 +25,13 @@
                                )
 
 process.maxEvents = cms.untracked.PSet(
-    input = cms.untracked.int32( 10 ))
+    input = cms.untracked.int32( 5 ))
+
+indir = '/afs/cern.ch/user/b/bfontana/CMSSW_11_2_0_pre5/src/23234.0_TTbar_14TeV+2026D49+TTbar_14TeV_TuneCP5_GenSimHLBeamSpot14+DigiTrigger+RecoGlobal+HARVESTGlobal'
+file_wildcard = 'step3.root'
+glob = glob.glob( os.path.join(indir, file_wildcard) )
+fNames = ['file:' + it for it in glob][:]
 
-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",
@@ -51,6 +56,7 @@
                                                      rcorr          = cms.vdouble( HGCalRecHit.__dict__['thicknessCorrection'][0:3] ),
                                                      weights        = HGCalRecHit.__dict__['layerWeights']
 )
+
 process.HeterogeneousHGCalHEFCellPositionsFiller = cms.ESProducer("HeterogeneousHGCalHEFCellPositionsFiller")
 process.HeterogeneousHGCalHEFRecHits = cms.EDProducer('HeterogeneousHGCalHEFRecHitProducer',
                                                       HGCHEFUncalibRecHitsTok = cms.InputTag('HGCalUncalibRecHit', 'HGCHEFUncalibRecHits'),
@@ -74,8 +80,8 @@
 process.HGCalRecHits = HGCalRecHit.clone()
 
 fNameOut = 'out'
-#process.task = cms.Task( process.HeterogeneousHGCalHEFCellPositionsFiller, process.HeterogeneousHGCalHEFRecHits )
-process.task = cms.Task( process.HGCalRecHits, process.HeterogeneousHGCalHEFRecHits )
+process.task = cms.Task( process.HeterogeneousHGCalHEFCellPositionsFiller, process.HeterogeneousHGCalHEFRecHits )
+#process.task = cms.Task( process.HGCalRecHits, process.HeterogeneousHGCalHEFRecHits )
 process.path = cms.Path( process.task )
 
 process.out = cms.OutputModule("PoolOutputModule",

UserCode/RecHitsValidator/python/HeterogeneousHGCalRecHitsValidator_cfg.py

@@ -20,17 +20,23 @@
 from Configuration.AlCa.GlobalTag import GlobalTag
 process.GlobalTag = GlobalTag(process.GlobalTag, 'auto:phase2_realistic', '')
 
-process.maxEvents = cms.untracked.PSet( input = cms.untracked.int32(-1) )
+process.maxEvents = cms.untracked.PSet( input = cms.untracked.int32(10) )
 
 #TFileService
+dirName = '/eos/user/b/bfontana/Samples/'
 fileName = 'validation.root'
 process.TFileService = cms.Service("TFileService", 
-                                   fileName = cms.string(fileName),
+                                   fileName = cms.string( os.path.join(dirName,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']
+fNames = ['file:/afs/cern.ch/user/b/bfontana/CMSSW_11_2_0_pre5/src/23234.0_TTbar_14TeV+2026D49+TTbar_14TeV_TuneCP5_GenSimHLBeamSpot14+DigiTrigger+RecoGlobal+HARVESTGlobal/step3.root']
+#indir = '/eos/cms/store/group/dpg_hgcal/comm_hgcal/bfontana/GPUScintillator/'
+#file_wildcard = 'step3_0.root'
+#glob = glob.glob( os.path.join(indir, file_wildcard) )
+#fNames = ['file:' + it for it in glob][:]
+
 keep = 'keep *'
 drop = 'drop CSCDetIdCSCALCTPreTriggerDigiMuonDigiCollection_simCscTriggerPrimitiveDigis__HLT'
 process.source = cms.Source("PoolSource",
@@ -85,22 +91,12 @@
                                 cpuRecHitsHSiToken = cms.InputTag('HGCalRecHits', 'HGCHEFRecHits'),
                                 gpuRecHitsHSiToken = cms.InputTag('HeterogeneousHGCalHEFRecHits','HeterogeneousHGCalHEFRecHits'),
                                 cpuRecHitsHSciToken = cms.InputTag('HGCalRecHits', 'HGCHEBRecHits'),
-                                gpuRecHitsHSciToken = cms.InputTag('HeterogeneousHGCalHEBRecHits','HeterogeneousHGCalHEBRecHits') )
-"""
-process.valid = cms.EDAnalyzer( "HeterogeneousHGCalRecHitsValidator",
-                                cpuRecHitsEEToken = cms.InputTag('HGCalRecHits', 'HGCHEFRecHits'),
-                                gpuRecHitsEEToken = cms.InputTag('HeterogeneousHGCalHEFRecHits','HeterogeneousHGCalHEFRecHits'),
-                                cpuRecHitsHSiToken = cms.InputTag('HGCalRecHits', 'HGCHEFRecHits'),
-                                gpuRecHitsHSiToken = cms.InputTag('HeterogeneousHGCalHEFRecHits','HeterogeneousHGCalHEFRecHits'),
-                                cpuRecHitsHSciToken = cms.InputTag('HGCalRecHits', 'HGCHEFRecHits'),
-                                gpuRecHitsHSciToken = cms.InputTag('HeterogeneousHGCalHEFRecHits','HeterogeneousHGCalHEFRecHits') 
+                                gpuRecHitsHSciToken = cms.InputTag('HeterogeneousHGCalHEBRecHits','HeterogeneousHGCalHEBRecHits')
 )
-"""
 
 process.recHitsTask = cms.Task( process.HGCalRecHits, process.HeterogeneousHGCalEERecHits, process.HeterogeneousHGCalHEFRecHits, process.HeterogeneousHGCalHEBRecHits )
-#process.recHitsTask = cms.Task( process.HGCalRecHits, process.HeterogeneousHGCalHEFRecHits )
 process.path = cms.Path( process.valid, process.recHitsTask )
 
 process.out = cms.OutputModule( "PoolOutputModule", 
-                                fileName = cms.untracked.string('out.root') )
+                                fileName = cms.untracked.string( os.path.join(dirName, 'out.root') ) )
 process.outpath = cms.EndPath(process.out)