Skip to content

Commit

Permalink
kernel maanger for ESProducer
Browse files Browse the repository at this point in the history
  • Loading branch information
bfonta committed Sep 7, 2020
1 parent 83b3702 commit 00683b6
Show file tree
Hide file tree
Showing 19 changed files with 219 additions and 122 deletions.
2 changes: 1 addition & 1 deletion .clang-format
@@ -1,4 +1,4 @@
---
macs -nw 2 ---
Language: Cpp
BasedOnStyle: Google
ColumnLimit: 120
Expand Down
29 changes: 18 additions & 11 deletions RecoLocalCalo/HGCalESProducers/BuildFile.xml
@@ -1,13 +1,20 @@
<use name="FWCore/MessageLogger"/>
<use name="FWCore/Framework"/>
<use name="FWCore/Utilities"/>
<use name="FWCore/ParameterSet"/>
<use name="RecoLocalCalo/HGCalRecAlgos"/>
<use name="RecoLocalCalo/HGCalRecProducers"/>
<use name="Geometry/HGCalGeometry"/>
<use name="Geometry/HGCalCommonData"/>
<use name="CondFormats/DataRecord"/>
<use name="clhep"/>
<use name="cuda"/>
<use name="FWCore/MessageLogger"/>
<use name="FWCore/Framework"/>
<use name="FWCore/Utilities"/>
<use name="FWCore/ParameterSet"/>
<use name="FWCore/ServiceRegistry"/>
<use name="Geometry/HGCalGeometry"/>
<use mame="Geometry/HGCalCommonData"/>
<use name="CondFormats/DataRecord"/>
<use name="CondFormats/HGCalObjects"/>
<use name="DataFormats/HGCRecHit"/>
<use name="DataFormats/ForwardDetId"/>
<use name="CommonTools/UtilAlgos"/>
<use name="RecoLocalCalo/HGCalRecAlgos"/>
<use name="HeterogeneousCore/CUDACore"/>
<use name="HeterogeneousCore/CUDAUtilities"/>
<use name="clhep"/>
<export>
<lib name="1"/>
<lib name="1"/>
</export>
12 changes: 9 additions & 3 deletions RecoLocalCalo/HGCalESProducers/plugins/BuildFile.xml
@@ -1,15 +1,21 @@
<use name="cuda"/>
<use name="HeterogeneousCore/CUDACore"/>
<use name="HeterogeneousCore/CUDAUtilities"/>
<use name="Geometry/CaloGeometry"/>
<use name="FWCore/Framework"/>
<use name="FWCore/Utilities"/>
<use name="FWCore/ParameterSet"/>
<use name="RecoLocalCalo/HGCalRecAlgos"/>
<use name="RecoLocalCalo/HGCalRecProducers"/>
<use name="DataFormats/HGCDigi"/>
<use name="DataFormats/HGCRecHit"/>
<use name="DataFormats/HcalDetId"/>
<use name="Geometry/CaloGeometry"/>
<use name="Geometry/HGCalGeometry"/>
<use name="Geometry/HGCalCommonData"/>
<use name="RecoLocalCalo/HGCalRecAlgos"/>
<use name="RecoLocalCalo/HGCalESProducers"/>
<use name="FWCore/MessageLogger"/>
<use name="CondFormats/DataRecord"/>
<use name="CondFormats/HGCalObjects"/>
<use name="PhysicsTools/UtilAlgos"/>
<library file="*.cc *.cu" name="RecoLocalCaloHGCalESProducersPlugins">
<flags EDM_PLUGIN="1"/>
</library>
@@ -0,0 +1,72 @@
#include <cuda.h>
#include <cuda_runtime.h>
#include <inttypes.h>
#include "DataFormats/ForwardDetId/interface/HGCalDetId.h"
#include "RecoLocalCalo/HGCalESProducers/plugins/HGCalCellPositionsKernelImpl.cuh"

__global__
void fill_positions_from_detids(const hgcal_conditions::HeterogeneousHEFCellPositionsConditionsESProduct* conds)
{
unsigned int tid = blockDim.x * blockIdx.x + threadIdx.x;

for (unsigned int i = tid; i < conds->nelems_posmap; i += blockDim.x * gridDim.x)
{
HeterogeneousHGCSiliconDetId did(conds->posmap.detid[i]);
const float cU = static_cast<float>( did.cellU() );
const float cV = static_cast<float>( did.cellV() );
const float wU = static_cast<float>( did.waferU() );
const float wV = static_cast<float>( did.waferV() );
const float ncells = static_cast<float>( did.nCells() );
const int32_t layer = did.layer();

//based on `std::pair<float, float> HGCalDDDConstants::locateCell(const HGCSiliconDetId&, bool)
const float r_x2 = conds->posmap.waferSize + conds->posmap.sensorSeparation;
const float r = 0.5f * r_x2;
const float sqrt3 = __fsqrt_rn(3.f);
const float rsqrt3 = __frsqrt_rn(3.f); //rsqrt: 1 / sqrt
const float R = r_x2 * rsqrt3;
const float n2 = ncells / 2.f;
const float yoff_abs = rsqrt3 * r_x2;
const float yoff = (layer%2==1) ? yoff_abs : -1.f * yoff_abs; //CHANGE according to Sunanda's reply
float xpos = (-2.f * wU + wV) * r;
float ypos = yoff + (1.5f * wV * R);
const float R1 = __fdividef( conds->posmap.waferSize, 3.f * ncells );
const float r1_x2 = R1 * sqrt3;
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. CAREFUL WITH LATER DETECTOR ALIGNMENT!!!
conds->posmap.y[i] = ypos;
}
}

__global__
void print_positions_from_detids(const hgcal_conditions::HeterogeneousHEFCellPositionsConditionsESProduct* conds)
{
unsigned int tid = blockDim.x * blockIdx.x + threadIdx.x;

for (unsigned int i = tid; i < conds->nelems_posmap; i += blockDim.x * gridDim.x)
{
HeterogeneousHGCSiliconDetId did(conds->posmap.detid[i]);
const int32_t layer = did.layer();
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);
}
}

__global__
void test(const hgcal_conditions::HeterogeneousHEFCellPositionsConditionsESProduct* conds)
{
unsigned int tid = blockDim.x * blockIdx.x + threadIdx.x;

if(tid == 0)
{
printf("Nelems: %u\n", static_cast<unsigned>(conds->nelems_posmap));
for(unsigned i=0; i<10; ++i)
{
printf("%lf ", conds->posmap.z_per_layer[i]);
printf("%lf ", conds->posmap.x[i]);
printf("\n");
}
}
}
@@ -0,0 +1,19 @@
#ifndef RecoLocalCalo_HGCalESProducers_HGCalCellPositionsKernelImpl_cuh
#define RecoLocalCalo_HGCalESProducers_HGCalCellPositionsKernelImpl_cuh

#include <cuda.h>
#include <cuda_runtime.h>

#include "CUDADataFormats/HGCal/interface/HGCConditions.h"
#include "RecoLocalCalo/HGCalESProducers/plugins/KernelManagerHGCalCellPositions.h"

__global__
void fill_positions_from_detids(const hgcal_conditions::HeterogeneousHEFCellPositionsConditionsESProduct* conds);

__global__
void print_positions_from_detids(const hgcal_conditions::HeterogeneousHEFCellPositionsConditionsESProduct* conds);

__global__
void test(const hgcal_conditions::HeterogeneousHEFCellPositionsConditionsESProduct* conds);

#endif //RecoLocalCalo_HGCalESProducers_HGCalCellPositionsKernelImpl_cuh
Expand Up @@ -22,7 +22,7 @@ void HeterogeneousHGCalHEFCellPositionsFiller::set_conditions_()

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_->z_per_layer.resize( (ddd_->lastLayer(true) - ddd_->firstLayer() + 1) * 2 );
posmap_->numberCellsHexagon.reserve(upper_estimate_wafer_number);
posmap_->detid.reserve(upper_estimate_cell_number);
//set positons-related variables
Expand All @@ -36,8 +36,8 @@ void HeterogeneousHGCalHEFCellPositionsFiller::set_conditions_()

//store detids following a geometry ordering
for(int ilayer=1; ilayer<=posmap_->lastLayer; ++ilayer) {
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
posmap_->z_per_layer[ilayer - 1] = static_cast<float>( ddd_->waferZ(ilayer, true) ); //originally a double
posmap_->z_per_layer[ilayer + (ddd_->lastLayer(true) - ddd_->firstLayer())] = 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) {
Expand Down Expand Up @@ -79,13 +79,11 @@ std::unique_ptr<HeterogeneousHGCalHEFCellPositionsConditions> HeterogeneousHGCal

set_conditions_();

/*
HeterogeneousHGCalHEFCellPositionsConditions esproduct(posmap_);
d_conds = esproduct.getHeterogeneousConditionsESProductAsync( 0 ); //could use ctx.stream()?
KernelManagerHGCalRecHit kernel_manager;
KernelManagerHGCalCellPositions kernel_manager( posmap_->detid.size() );
kernel_manager.fill_positions(d_conds);
std::unique_ptr<HeterogeneousHGCalHEFCellPositionsConditionsESProduct> up(d_conds);
*/
//std::unique_ptr<HeterogeneousHGCalHEFCellPositionsConditionsESProduct> up(d_conds);

std::unique_ptr<HeterogeneousHGCalHEFCellPositionsConditions> up = std::make_unique<HeterogeneousHGCalHEFCellPositionsConditions>(posmap_);
return up;
Expand Down
Expand Up @@ -20,16 +20,13 @@
#include "FWCore/ParameterSet/interface/ConfigurationDescriptions.h"
#include "FWCore/ParameterSet/interface/ParameterSetDescription.h"

#include "RecoLocalCalo/HGCalRecAlgos/interface/RecHitTools.h"

#include "Geometry/HGCalGeometry/interface/HGCalGeometry.h"
#include "Geometry/HGCalCommonData/interface/HGCalDDDConstants.h"
#include "Geometry/HGCalCommonData/interface/HGCalWaferIndex.h"

#include "CondFormats/HGCalObjects/interface/HeterogeneousHGCalHEFCellPositionsConditions.h"
#include "RecoLocalCalo/HGCalRecProducers/plugins/HeterogeneousHGCalProducerMemoryWrapper.h"
#include "RecoLocalCalo/HGCalRecProducers/plugins/KernelManagerHGCalRecHit.h"
#include "CondFormats/DataRecord/interface/HeterogeneousHGCalHEFCellPositionsConditionsRecord.h"
#include "RecoLocalCalo/HGCalESProducers/plugins/KernelManagerHGCalCellPositions.h"

class HeterogeneousHGCalHEFCellPositionsFiller: public edm::ESProducer
{
Expand Down
@@ -0,0 +1,22 @@
#include <cuda.h>
#include <cuda_runtime.h>
#include <inttypes.h>
#include "RecoLocalCalo/HGCalESProducers/plugins/KernelManagerHGCalCellPositions.h"
#include "RecoLocalCalo/HGCalESProducers/plugins/HGCalCellPositionsKernelImpl.cuh"

KernelManagerHGCalCellPositions::KernelManagerHGCalCellPositions(const size_t& nelems)
{
::nb_celpos_ = (nelems + ::nt_celpos_.x - 1) / ::nt_celpos_.x;
}

void KernelManagerHGCalCellPositions::fill_positions(const hgcal_conditions::HeterogeneousHEFCellPositionsConditionsESProduct* d_conds)
{
fill_positions_from_detids<<<::nb_celpos_,::nt_celpos_>>>(d_conds);
cudaCheck( cudaGetLastError() );
}

void KernelManagerHGCalCellPositions::test_cell_positions(const hgcal_conditions::HeterogeneousHEFCellPositionsConditionsESProduct* d_conds)
{
test<<<::nb_celpos_,::nt_celpos_>>>(d_conds);
cudaCheck( cudaGetLastError() );
}
@@ -0,0 +1,37 @@
#ifndef RecoLocalCalo_HGCalESProducers_KernelManagerHGCalCellPositions_h
#define RecoLocalCalo_HGCalESProducers_KernelManagerHGCalCellPositions_h

#include "FWCore/Utilities/interface/Exception.h"
#include "HeterogeneousCore/CUDAUtilities/interface/MessageLogger.h"
#include "DataFormats/DetId/interface/DetId.h"
#include "HeterogeneousCore/CUDAUtilities/interface/cudaCompat.h"
#include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h"
#include "RecoLocalCalo/HGCalESProducers/plugins/HGCalCellPositionsKernelImpl.cuh"
#include "CUDADataFormats/HGCal/interface/HGCConditions.h"

#include <vector>
#include <algorithm> //std::swap
#include <variant>
#include <cuda.h>
#include <cuda_runtime.h>

/*
#ifdef __CUDA_ARCH__
extern __constant__ uint32_t calo_rechit_masks[];
#endif
*/

namespace { //kernel parameters
dim3 nb_celpos_;
constexpr dim3 nt_celpos_(256);
}

class KernelManagerHGCalCellPositions {
public:
KernelManagerHGCalCellPositions(const size_t&);

void fill_positions(const hgcal_conditions::HeterogeneousHEFCellPositionsConditionsESProduct*);
void test_cell_positions(const hgcal_conditions::HeterogeneousHEFCellPositionsConditionsESProduct*);
};

#endif //RecoLocalCalo_HGCalESProducers_KernelManagerHGCalCellPositions_h
6 changes: 4 additions & 2 deletions RecoLocalCalo/HGCalRecProducers/BuildFile.xml
@@ -1,18 +1,20 @@
<use name="cuda"/>
<use name="FWCore/Framework"/>
<use name="FWCore/Utilities"/>
<use name="FWCore/ParameterSet"/>
<use name="FWCore/ServiceRegistry"/>
<use name="FWCore/MessageLogger"/>
<use name="HeterogeneousCore/CUDACore"/>
<use name="HeterogeneousCore/CUDAUtilities"/>
<use name="DataFormats/HGCRecHit"/>
<use name="DataFormats/ForwardDetId"/>
<use name="CommonTools/UtilAlgos"/>
<use name="RecoLocalCalo/HGCalRecAlgos"/>
<use name="RecoLocalCalo/HGCalESProducers"/>
<use name="Geometry/HGCalGeometry"/>
<use name="Geometry/HGCalCommonData"/>
<use name="CondFormats/DataRecord"/>
<use name="CondFormats/HGCalObjects"/>
<use name="HeterogeneousCore/CUDACore"/>
<use name="HeterogeneousCore/CUDAUtilities"/>
<use name="clhep"/>
<export>
<lib name="1"/>
Expand Down
4 changes: 3 additions & 1 deletion RecoLocalCalo/HGCalRecProducers/plugins/BuildFile.xml
@@ -1,7 +1,6 @@
<use name="cuda"/>
<use name="HeterogeneousCore/CUDACore"/>
<use name="HeterogeneousCore/CUDAUtilities"/>
<use name="Geometry/CaloGeometry"/>
<use name="FWCore/Framework"/>
<use name="FWCore/Utilities"/>
<use name="FWCore/ParameterSet"/>
Expand All @@ -10,11 +9,14 @@
<use name="DataFormats/HcalDetId"/>
<use name="RecoLocalCalo/HGCalRecAlgos"/>
<use name="RecoLocalCalo/HGCalRecProducers"/>
<use name="RecoLocalCalo/HGCalESProducers"/>
<use name="FWCore/MessageLogger"/>
<use name="Geometry/CaloGeometry"/>
<use name="Geometry/HGCalGeometry"/>
<use name="Geometry/HGCalCommonData"/>
<use name="PhysicsTools/UtilAlgos"/>
<use name="CondFormats/DataRecord"/>
<use name="CondFormats/HGCalObjects"/>
<library file="*.cc *.cu" name="RecoLocalCaloHGCalRecProducersPlugins">
<flags EDM_PLUGIN="1"/>
</library>
Expand Up @@ -247,4 +247,4 @@ void HGCalLayerClusterProducer::produce(edm::Event& evt, const edm::EventSetup&
algo->reset();
}

#endif
#endif //__RecoLocalCalo_HGCRecProducers_HGCalLayerClusterProducer_H__
50 changes: 0 additions & 50 deletions RecoLocalCalo/HGCalRecProducers/plugins/HGCalRecHitKernelImpl.cu
Expand Up @@ -158,53 +158,3 @@ void heb_to_rechit(HGCRecHitSoA dst_soa, HGCUncalibratedRecHitSoA src_soa, const
make_rechit_scintillator(i, dst_soa, src_soa, weight, sigmaNoiseGeV);
}
}

__global__
void fill_positions_from_detids(const hgcal_conditions::HeterogeneousHEFCellPositionsConditionsESProduct* conds)
{
unsigned int tid = blockDim.x * blockIdx.x + threadIdx.x;

for (unsigned int i = tid; i < conds->nelems_posmap; i += blockDim.x * gridDim.x)
{
HeterogeneousHGCSiliconDetId did(conds->posmap.detid[i]);
const float cU = static_cast<float>( did.cellU() );
const float cV = static_cast<float>( did.cellV() );
const float wU = static_cast<float>( did.waferU() );
const float wV = static_cast<float>( did.waferV() );
const float ncells = static_cast<float>( did.nCells() );
const int32_t layer = did.layer();

//based on `std::pair<float, float> HGCalDDDConstants::locateCell(const HGCSiliconDetId&, bool)
const float r_x2 = conds->posmap.waferSize + conds->posmap.sensorSeparation;
const float r = 0.5f * r_x2;
const float sqrt3 = __fsqrt_rn(3.f);
const float rsqrt3 = __frsqrt_rn(3.f); //rsqrt: 1 / sqrt
const float R = r_x2 * rsqrt3;
const float n2 = ncells / 2.f;
const float yoff_abs = rsqrt3 * r_x2;
const float yoff = (layer%2==1) ? yoff_abs : -1.f * yoff_abs; //CHANGE according to Sunanda's reply
float xpos = (-2.f * wU + wV) * r;
float ypos = yoff + (1.5f * wV * R);
const float R1 = __fdividef( conds->posmap.waferSize, 3.f * ncells );
const float r1_x2 = R1 * sqrt3;
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. CAREFUL WITH LATER DETECTOR ALIGNMENT!!!
conds->posmap.y[i] = ypos;
}
}

__global__
void print_positions_from_detids(const hgcal_conditions::HeterogeneousHEFCellPositionsConditionsESProduct* conds)
{
unsigned int tid = blockDim.x * blockIdx.x + threadIdx.x;

for (unsigned int i = tid; i < conds->nelems_posmap; i += blockDim.x * gridDim.x)
{
HeterogeneousHGCSiliconDetId did(conds->posmap.detid[i]);
const int32_t layer = did.layer();
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);
}
}
Expand Up @@ -7,7 +7,6 @@
#include "CUDADataFormats/HGCal/interface/HGCUncalibratedRecHitSoA.h"
#include "CUDADataFormats/HGCal/interface/HGCRecHitSoA.h"
#include "CUDADataFormats/HGCal/interface/HGCUncalibratedRecHitsToRecHitsConstants.h"
#include "CUDADataFormats/HGCal/interface/HGCConditions.h"

#include "RecoLocalCalo/HGCalRecProducers/plugins/KernelManagerHGCalRecHit.h"

Expand All @@ -28,11 +27,5 @@ void hef_to_rechit(HGCRecHitSoA dst_soa, HGCUncalibratedRecHitSoA src_soa, HGChe

__global__
void heb_to_rechit(HGCRecHitSoA dst_soa, HGCUncalibratedRecHitSoA src_soa, HGChebUncalibratedRecHitConstantData cdata, int length);

__global__
void fill_positions_from_detids(const hgcal_conditions::HeterogeneousHEFCellPositionsConditionsESProduct*);

__global__
void print_positions_from_detids(const hgcal_conditions::HeterogeneousHEFCellPositionsConditionsESProduct*);

#endif //RecoLocalCalo_HGCalRecProducers_HGCalRecHitKernelImpl_cuh

0 comments on commit 00683b6

Please sign in to comment.