Warning, /RecoLocalCalo/EcalRecProducers/plugins/EcalRecHitBuilderKernels.cu is written in an unsupported language. File is not indexed.
0001 #include <cuda.h>
0002
0003 #include "CUDADataFormats/EcalRecHitSoA/interface/EcalRecHit.h"
0004 #include "CUDADataFormats/EcalRecHitSoA/interface/EcalUncalibratedRecHit.h"
0005
0006 #include "EcalRecHitBuilderKernels.h"
0007 #include "KernelHelpers.h"
0008
0009 namespace ecal {
0010 namespace rechit {
0011
0012 // uncalibrecHit flags
0013 enum UncalibRecHitFlags {
0014 kGood = -1, // channel is good (mutually exclusive with other states) setFlagBit(kGood) reset flags_ to zero
0015 kPoorReco, // channel has been badly reconstructed (e.g. bad shape, bad chi2 etc.)
0016 kSaturated, // saturated channel
0017 kOutOfTime, // channel out of time
0018 kLeadingEdgeRecovered, // saturated channel: energy estimated from the leading edge before saturation
0019 kHasSwitchToGain6, // at least one data frame is in G6
0020 kHasSwitchToGain1 // at least one data frame is in G1
0021 };
0022
0023 // recHit flags
0024 enum RecHitFlags {
0025 RecHitFlags_kGood = 0, // channel ok, the energy and time measurement are reliable
0026 RecHitFlags_kPoorReco, // the energy is available from the UncalibRecHit, but approximate (bad shape, large chi2)
0027 RecHitFlags_kOutOfTime, // the energy is available from the UncalibRecHit (sync reco), but the event is out of time
0028 RecHitFlags_kFaultyHardware, // The energy is available from the UncalibRecHit, channel is faulty at some hardware level (e.g. noisy)
0029 RecHitFlags_kNoisy, // the channel is very noisy
0030 RecHitFlags_kPoorCalib, // the energy is available from the UncalibRecHit, but the calibration of the channel is poor
0031 RecHitFlags_kSaturated, // saturated channel (recovery not tried)
0032 RecHitFlags_kLeadingEdgeRecovered, // saturated channel: energy estimated from the leading edge before saturation
0033 RecHitFlags_kNeighboursRecovered, // saturated/isolated dead: energy estimated from neighbours
0034 RecHitFlags_kTowerRecovered, // channel in TT with no data link, info retrieved from Trigger Primitive
0035 RecHitFlags_kDead, // channel is dead and any recovery fails
0036 RecHitFlags_kKilled, // MC only flag: the channel is killed in the real detector
0037 RecHitFlags_kTPSaturated, // the channel is in a region with saturated TP
0038 RecHitFlags_kL1SpikeFlag, // the channel is in a region with TP with sFGVB = 0
0039 RecHitFlags_kWeird, // the signal is believed to originate from an anomalous deposit (spike)
0040 RecHitFlags_kDiWeird, // the signal is anomalous, and neighbors another anomalous signal
0041 RecHitFlags_kHasSwitchToGain6, // at least one data frame is in G6
0042 RecHitFlags_kHasSwitchToGain1, // at least one data frame is in G1
0043 //
0044 RecHitFlags_kUnknown // to ease the interface with functions returning flags.
0045 };
0046
0047 // status code
0048 enum EcalChannelStatusCode_Code {
0049 kOk = 0,
0050 kDAC,
0051 kNoLaser,
0052 kNoisy,
0053 kNNoisy,
0054 kNNNoisy,
0055 kNNNNoisy,
0056 kNNNNNoisy,
0057 kFixedG6,
0058 kFixedG1,
0059 kFixedG0,
0060 kNonRespondingIsolated,
0061 kDeadVFE,
0062 kDeadFE,
0063 kNoDataNoTP
0064 };
0065
0066 __global__ void kernel_create_ecal_rehit(
0067 // configuration
0068 int const* ChannelStatusToBeExcluded,
0069 uint32_t ChannelStatusToBeExcludedSize,
0070 bool const killDeadChannels,
0071 bool const recoverEBIsolatedChannels,
0072 bool const recoverEEIsolatedChannels,
0073 bool const recoverEBVFE,
0074 bool const recoverEEVFE,
0075 bool const recoverEBFE,
0076 bool const recoverEEFE,
0077 float const EBLaserMIN,
0078 float const EELaserMIN,
0079 float const EBLaserMAX,
0080 float const EELaserMAX,
0081 // for flags setting
0082 int const* expanded_v_DB_reco_flags, // FIXME AM: to be checked
0083 uint32_t const* expanded_Sizes_v_DB_reco_flags,
0084 uint32_t const* expanded_flagbit_v_DB_reco_flags,
0085 uint32_t expanded_v_DB_reco_flagsSize,
0086 uint32_t flagmask,
0087 // conditions
0088 float const* adc2gev,
0089 float const* intercalib,
0090 uint16_t const* status,
0091 float const* apdpnrefs,
0092 float const* alphas,
0093 // input for transparency corrections
0094 float const* p1,
0095 float const* p2,
0096 float const* p3,
0097 edm::TimeValue_t const* t1,
0098 edm::TimeValue_t const* t2,
0099 edm::TimeValue_t const* t3,
0100 // input for linear corrections
0101 float const* lp1,
0102 float const* lp2,
0103 float const* lp3,
0104 edm::TimeValue_t const* lt1,
0105 edm::TimeValue_t const* lt2,
0106 edm::TimeValue_t const* lt3,
0107 // time, used for time dependent corrections
0108 edm::TimeValue_t const event_time,
0109 // input
0110 uint32_t const* did_eb,
0111 uint32_t const* did_ee,
0112 ::ecal::reco::StorageScalarType const* amplitude_eb, // in adc counts
0113 ::ecal::reco::StorageScalarType const* amplitude_ee, // in adc counts
0114 ::ecal::reco::StorageScalarType const* time_eb,
0115 ::ecal::reco::StorageScalarType const* time_ee,
0116 ::ecal::reco::StorageScalarType const* chi2_eb,
0117 ::ecal::reco::StorageScalarType const* chi2_ee,
0118 uint32_t const* flags_eb,
0119 uint32_t const* flags_ee,
0120 // output
0121 uint32_t* didEB,
0122 uint32_t* didEE,
0123 ::ecal::reco::StorageScalarType* energyEB, // in energy [GeV]
0124 ::ecal::reco::StorageScalarType* energyEE, // in energy [GeV]
0125 ::ecal::reco::StorageScalarType* timeEB,
0126 ::ecal::reco::StorageScalarType* timeEE,
0127 ::ecal::reco::StorageScalarType* chi2EB,
0128 ::ecal::reco::StorageScalarType* chi2EE,
0129 uint32_t* flagBitsEB,
0130 uint32_t* flagBitsEE,
0131 uint32_t* extraEB,
0132 uint32_t* extraEE,
0133 // other
0134 int const nchannels,
0135 uint32_t const nChannelsBarrel,
0136 uint32_t const offsetForHashes) {
0137 //
0138 // NB: energy "type_wrapper<reco::StorageScalarType, L>::type" most likely std::vector<float>
0139 //
0140
0141 for (int ch = threadIdx.x + blockDim.x * blockIdx.x; ch < nchannels; ch += blockDim.x * gridDim.x) {
0142 bool isEndcap = (ch >= nChannelsBarrel);
0143
0144 int const inputCh = isEndcap ? ch - nChannelsBarrel : ch;
0145
0146 uint32_t const* didCh = isEndcap ? did_ee : did_eb;
0147
0148 // arrange to access the right ptrs
0149 #define ARRANGE(var) auto* var = isEndcap ? var##EE : var##EB
0150 ARRANGE(did);
0151 ARRANGE(energy);
0152 ARRANGE(chi2);
0153 ARRANGE(flagBits);
0154 ARRANGE(extra);
0155 #undef ARRANGE
0156
0157 // only two values, EB or EE
0158 // AM : FIXME : why not using "isBarrel" ? isBarrel ? adc2gev[0] : adc2gev[1]
0159 float adc2gev_to_use = isEndcap ? adc2gev[1] // ee
0160 : adc2gev[0]; // eb
0161
0162 // first EB and then EE
0163
0164 ::ecal::reco::StorageScalarType const* amplitude = isEndcap ? amplitude_ee : amplitude_eb;
0165
0166 ::ecal::reco::StorageScalarType const* chi2_in = isEndcap ? chi2_ee : chi2_eb;
0167
0168 uint32_t const* flags_in = isEndcap ? flags_ee : flags_eb;
0169
0170 // simple copy
0171 did[inputCh] = didCh[inputCh];
0172
0173 auto const did_to_use = DetId{didCh[inputCh]};
0174
0175 auto const isBarrel = did_to_use.subdetId() == EcalBarrel;
0176 auto const hashedId = isBarrel ? ecal::reconstruction::hashedIndexEB(did_to_use.rawId())
0177 : offsetForHashes + ecal::reconstruction::hashedIndexEE(did_to_use.rawId());
0178
0179 float const intercalib_to_use = intercalib[hashedId];
0180
0181 // get laser coefficient
0182 float lasercalib = 1.;
0183
0184 //
0185 // AM: ideas
0186 //
0187 // One possibility is to create the map of laser corrections once on CPU
0188 // for all crystals and push them on GPU.
0189 // Then only if the LS is different, update the laser correction
0190 // The variation within a LS is not worth pursuing (<< 0.1% !!)
0191 // and below the precision we can claim on the laser corrections (right?).
0192 // This will save quite some time (also for the CPU version?)
0193 //
0194
0195 int iLM = 1;
0196
0197 if (isBarrel) {
0198 iLM = ecal::reconstruction::laser_monitoring_region_EB(did_to_use.rawId());
0199 } else {
0200 iLM = ecal::reconstruction::laser_monitoring_region_EE(did_to_use.rawId());
0201 }
0202
0203 long long t_i = 0, t_f = 0;
0204 float p_i = 0, p_f = 0;
0205 long long lt_i = 0, lt_f = 0;
0206 float lp_i = 0, lp_f = 0;
0207
0208 // laser
0209 if (event_time >= t1[iLM - 1] && event_time < t2[iLM - 1]) {
0210 t_i = t1[iLM - 1];
0211 t_f = t2[iLM - 1];
0212 p_i = p1[hashedId];
0213 p_f = p2[hashedId];
0214 } else if (event_time >= t2[iLM - 1] && event_time <= t3[iLM - 1]) {
0215 t_i = t2[iLM - 1];
0216 t_f = t3[iLM - 1];
0217 p_i = p2[hashedId];
0218 p_f = p3[hashedId];
0219 } else if (event_time < t1[iLM - 1]) {
0220 t_i = t1[iLM - 1];
0221 t_f = t2[iLM - 1];
0222 p_i = p1[hashedId];
0223 p_f = p2[hashedId];
0224
0225 } else if (event_time > t3[iLM - 1]) {
0226 t_i = t2[iLM - 1];
0227 t_f = t3[iLM - 1];
0228 p_i = p2[hashedId];
0229 p_f = p3[hashedId];
0230 }
0231
0232 // linear corrections
0233 if (event_time >= lt1[iLM - 1] && event_time < lt2[iLM - 1]) {
0234 lt_i = lt1[iLM - 1];
0235 lt_f = lt2[iLM - 1];
0236 lp_i = lp1[hashedId];
0237 lp_f = lp2[hashedId];
0238 } else if (event_time >= lt2[iLM - 1] && event_time <= lt3[iLM - 1]) {
0239 lt_i = lt2[iLM - 1];
0240 lt_f = lt3[iLM - 1];
0241 lp_i = lp2[hashedId];
0242 lp_f = lp3[hashedId];
0243 } else if (event_time < lt1[iLM - 1]) {
0244 lt_i = lt1[iLM - 1];
0245 lt_f = lt2[iLM - 1];
0246 lp_i = lp1[hashedId];
0247 lp_f = lp2[hashedId];
0248
0249 } else if (event_time > lt3[iLM - 1]) {
0250 lt_i = lt2[iLM - 1];
0251 lt_f = lt3[iLM - 1];
0252 lp_i = lp2[hashedId];
0253 lp_f = lp3[hashedId];
0254 }
0255
0256 // apdpnref and alpha
0257 float apdpnref = apdpnrefs[hashedId];
0258 float alpha = alphas[hashedId];
0259
0260 // now calculate transparency correction
0261 if (apdpnref != 0 && (t_i - t_f) != 0 && (lt_i - lt_f) != 0) {
0262 long long tt = event_time; // never subtract two unsigned!
0263 float interpolatedLaserResponse =
0264 p_i / apdpnref + float(tt - t_i) * (p_f - p_i) / (apdpnref * float(t_f - t_i));
0265
0266 float interpolatedLinearResponse =
0267 lp_i / apdpnref + float(tt - lt_i) * (lp_f - lp_i) / (apdpnref * float(lt_f - lt_i)); // FIXED BY FC
0268
0269 if (interpolatedLinearResponse > 2.f || interpolatedLinearResponse < 0.1f) {
0270 interpolatedLinearResponse = 1.f;
0271 }
0272 if (interpolatedLaserResponse <= 0.) {
0273 // AM : how the heck is it possible?
0274 // interpolatedLaserResponse = 0.0001;
0275 lasercalib = 1.;
0276
0277 } else {
0278 float interpolatedTransparencyResponse = interpolatedLaserResponse / interpolatedLinearResponse;
0279
0280 // ... and now this:
0281 lasercalib = 1.f / (std::pow(interpolatedTransparencyResponse, alpha) * interpolatedLinearResponse);
0282 }
0283 }
0284
0285 //
0286 // Check for channels to be excluded from reconstruction
0287 //
0288 // Default energy not to be updated if "ChannelStatusToBeExcluded"
0289 // Exploited later by the module "EcalRecHitConvertGPU2CPUFormat"
0290 energy[inputCh] = -1; //un-physical default
0291
0292 // truncate the chi2
0293 if (chi2_in[inputCh] > 64)
0294 chi2[inputCh] = 64;
0295 else
0296 chi2[inputCh] = chi2_in[inputCh];
0297
0298 // default values for the flags
0299 flagBits[inputCh] = 0;
0300 extra[inputCh] = 0;
0301
0302 static const int chStatusMask = 0x1f;
0303 // ChannelStatusToBeExcluded is a "int" then I put "dbstatus" to be the same
0304 int dbstatus = EcalChannelStatusCode_Code((status[hashedId]) & chStatusMask);
0305 if (ChannelStatusToBeExcludedSize != 0) {
0306 bool skip_this_channel = false;
0307 for (int ich_to_check = 0; ich_to_check < ChannelStatusToBeExcludedSize; ich_to_check++) {
0308 if (ChannelStatusToBeExcluded[ich_to_check] == dbstatus) {
0309 skip_this_channel = true;
0310 break;
0311 }
0312 }
0313 if (skip_this_channel) {
0314 // skip this channel
0315 continue;
0316 }
0317 }
0318
0319 // Take our association map of dbstatuses-> recHit flagbits and return the apporpriate flagbit word
0320
0321 //
0322 // AM: get the smaller "flagbit_counter" with match
0323 //
0324
0325 uint32_t temporary_flagBits = 0;
0326
0327 int iterator_flags = 0;
0328 bool need_to_exit = false;
0329 int flagbit_counter = 0;
0330 while (!need_to_exit) {
0331 iterator_flags = 0;
0332 for (unsigned int i = 0; i != expanded_v_DB_reco_flagsSize; ++i) {
0333 // check the correct "flagbit"
0334 if (expanded_flagbit_v_DB_reco_flags[i] == flagbit_counter) {
0335 for (unsigned int j = 0; j < expanded_Sizes_v_DB_reco_flags[i]; j++) {
0336 if (expanded_v_DB_reco_flags[iterator_flags] == dbstatus) {
0337 temporary_flagBits = 0x1 << expanded_flagbit_v_DB_reco_flags[i];
0338 need_to_exit = true;
0339 break; // also from the big loop!!!
0340 }
0341 iterator_flags++;
0342 }
0343 } else {
0344 // if not, got to the next bunch directly
0345 iterator_flags += expanded_Sizes_v_DB_reco_flags[i];
0346 }
0347
0348 if (need_to_exit) {
0349 break;
0350 }
0351 }
0352 flagbit_counter += 1;
0353 }
0354
0355 flagBits[inputCh] = temporary_flagBits;
0356
0357 if ((flagmask & temporary_flagBits) && killDeadChannels) {
0358 // skip this channel
0359 continue;
0360 }
0361
0362 //
0363 // multiply the adc counts with factors to get GeV
0364 //
0365
0366 // energy[ch] = amplitude[inputCh] * adc2gev_to_use * intercalib_to_use ;
0367 energy[inputCh] = amplitude[inputCh] * adc2gev_to_use * intercalib_to_use * lasercalib;
0368
0369 // Time is not saved so far, FIXME
0370 // time[ch] = time_in[inputCh];
0371
0372 // NB: calculate the "flagBits extra" --> not really "flags", but actually an encoded version of energy uncertainty, time unc., ...
0373
0374 //
0375 // extra packing ...
0376 //
0377
0378 uint32_t offset;
0379 uint32_t width;
0380 uint32_t value;
0381
0382 float chi2_temp = chi2[inputCh];
0383 if (chi2_temp > 64)
0384 chi2_temp = 64;
0385 // use 7 bits
0386 uint32_t rawChi2 = lround(chi2_temp / 64. * ((1 << 7) - 1));
0387
0388 offset = 0;
0389 width = 7;
0390 value = 0;
0391
0392 uint32_t mask = ((1 << width) - 1) << offset;
0393 value &= ~mask;
0394 value |= (rawChi2 & ((1U << width) - 1)) << offset;
0395
0396 // rawEnergy is actually "error" !!!
0397 uint32_t rawEnergy = 0;
0398
0399 // AM: FIXME: this is not propagated currently to the uncalibrecHit collection SOA
0400 // if you want to store this in "extra", we need first to add it to the uncalibrecHit results
0401 // then it will be something like the following
0402 // amplitudeError[inputCh] * adc2gev_to_use * intercalib_to_use * lasercalib
0403 //
0404 //
0405
0406 float amplitudeError_ch = 0.; // amplitudeError[ch];
0407
0408 if (amplitudeError_ch > 0.001) {
0409 static constexpr float p10[] = {1.e-2f, 1.e-1f, 1.f, 1.e1f, 1.e2f, 1.e3f, 1.e4f, 1.e5f, 1.e6f};
0410 int b = amplitudeError_ch < p10[4] ? 0 : 5;
0411 for (; b < 9; ++b)
0412 if (amplitudeError_ch < p10[b])
0413 break;
0414
0415 uint16_t exponent = b;
0416
0417 static constexpr float ip10[] = {1.e5f, 1.e4f, 1.e3f, 1.e2f, 1.e1f, 1.e0f, 1.e-1f, 1.e-2f, 1.e-3f, 1.e-4};
0418 uint16_t significand = lround(amplitudeError_ch * ip10[exponent]);
0419 // use 13 bits (3 exponent, 10 significand)
0420 rawEnergy = exponent << 10 | significand;
0421 }
0422
0423 offset = 8;
0424 width = 13;
0425 // value from last change, ok
0426
0427 mask = ((1 << width) - 1) << offset;
0428 value &= ~mask;
0429 value |= (rawEnergy & ((1U << width) - 1)) << offset;
0430
0431 uint32_t jitterErrorBits = 0;
0432 jitterErrorBits = jitterErrorBits & 0xFF;
0433
0434 offset = 24;
0435 width = 8;
0436 // value from last change, ok
0437
0438 mask = ((1 << width) - 1) << offset;
0439 value &= ~mask;
0440 value |= (jitterErrorBits & ((1U << width) - 1)) << offset;
0441
0442 //
0443 // now finally set "extra[ch]"
0444 //
0445 extra[inputCh] = value;
0446
0447 //
0448 // additional flags setting
0449 //
0450 // using correctly the flags as calculated at the UncalibRecHit stage
0451 //
0452 // Now fill flags
0453
0454 bool good = true;
0455
0456 if (flags_in[inputCh] & (0x1 << (UncalibRecHitFlags::kLeadingEdgeRecovered))) {
0457 flagBits[inputCh] |= (0x1 << (RecHitFlags::RecHitFlags_kLeadingEdgeRecovered));
0458 good = false;
0459 }
0460
0461 if (flags_in[inputCh] & (0x1 << (UncalibRecHitFlags::kSaturated))) {
0462 // leading edge recovery failed - still keep the information
0463 // about the saturation and do not flag as dead
0464 flagBits[inputCh] |= (0x1 << (RecHitFlags::RecHitFlags_kSaturated));
0465 good = false;
0466 }
0467
0468 //
0469 // AM: why do we have two tests one after the other checking almost the same thing???
0470 // Please clean up the code, ... also the original one!
0471 //
0472 // uncalibRH.isSaturated() --->
0473 //
0474 // bool EcalUncalibratedRecHit::isSaturated() const {
0475 // return EcalUncalibratedRecHit::checkFlag(kSaturated);
0476 // }
0477 //
0478 //
0479
0480 if (flags_in[inputCh] & (0x1 << (UncalibRecHitFlags::kSaturated))) {
0481 flagBits[inputCh] |= (0x1 << (RecHitFlags::RecHitFlags_kSaturated));
0482 good = false;
0483 }
0484
0485 if (flags_in[inputCh] & (0x1 << (UncalibRecHitFlags::kOutOfTime))) {
0486 flagBits[inputCh] |= (0x1 << (RecHitFlags::RecHitFlags_kOutOfTime));
0487 good = false;
0488 }
0489 if (flags_in[inputCh] & (0x1 << (UncalibRecHitFlags::kPoorReco))) {
0490 flagBits[inputCh] |= (0x1 << (RecHitFlags::RecHitFlags_kPoorReco));
0491 good = false;
0492 }
0493 if (flags_in[inputCh] & (0x1 << (UncalibRecHitFlags::kHasSwitchToGain6))) {
0494 flagBits[inputCh] |= (0x1 << (RecHitFlags::RecHitFlags_kHasSwitchToGain6));
0495 }
0496 if (flags_in[inputCh] & (0x1 << (UncalibRecHitFlags::kHasSwitchToGain1))) {
0497 flagBits[inputCh] |= (0x1 << (RecHitFlags::RecHitFlags_kHasSwitchToGain1));
0498 }
0499
0500 if (good) {
0501 flagBits[inputCh] |= (0x1 << (RecHitFlags::RecHitFlags_kGood));
0502 }
0503
0504 if ((isBarrel && (lasercalib < EBLaserMIN || lasercalib > EBLaserMAX)) ||
0505 (!isBarrel && (lasercalib < EELaserMIN || lasercalib > EELaserMAX))) {
0506 flagBits[inputCh] |= (0x1 << (RecHitFlags::RecHitFlags_kPoorCalib));
0507 }
0508
0509 // recover, killing, and other stuff
0510
0511 //
0512 // Structure:
0513 // EB
0514 // EE
0515 //
0516 //
0517 // - single MVA
0518 // - democratic sharing
0519 // - kill all the other cases
0520 //
0521
0522 bool is_Single = false;
0523 bool is_FE = false;
0524 bool is_VFE = false;
0525
0526 bool is_recoverable = false; // DetIdToBeRecovered
0527
0528 if (dbstatus == 10 || dbstatus == 11 || dbstatus == 12) {
0529 is_recoverable = true;
0530 }
0531
0532 if (is_recoverable) {
0533 if (dbstatus == EcalChannelStatusCode_Code::kDeadVFE) {
0534 is_VFE = true;
0535 } else if (dbstatus == EcalChannelStatusCode_Code::kDeadVFE) {
0536 is_FE = true;
0537 } else {
0538 is_Single = true;
0539 }
0540
0541 // EB
0542 if (isBarrel) {
0543 if (is_Single || is_FE || is_VFE) {
0544 // single MVA
0545 if (is_Single && (recoverEBIsolatedChannels || !killDeadChannels)) {
0546 }
0547 // decmocratic sharing
0548 else if (is_FE && (recoverEBFE || !killDeadChannels)) {
0549 }
0550 // kill all the other cases
0551 else {
0552 energy[inputCh] = 0.; // Need to set also the flags ...
0553 }
0554 }
0555 }
0556 // EE
0557 else {
0558 if (is_Single || is_FE || is_VFE) {
0559 // single MVA
0560 if (is_Single && (recoverEBIsolatedChannels || !killDeadChannels)) {
0561 }
0562 // decmocratic sharing
0563 else if (is_FE && (recoverEBFE || !killDeadChannels)) {
0564 //
0565 // Code is definitely too long ...
0566 //
0567
0568 }
0569 // kill all the other cases
0570 else {
0571 energy[inputCh] = 0.; // Need to set also the flags ...
0572 }
0573 }
0574 }
0575 }
0576
0577 } // end channel
0578 }
0579
0580 // host version, to be called by the plugin
0581 void create_ecal_rehit(EventInputDataGPU const& eventInputGPU,
0582 EventOutputDataGPU& eventOutputGPU,
0583 // eventDataForScratchGPU_,
0584 ConditionsProducts const& conditions,
0585 ConfigurationParameters const& configParameters,
0586 uint32_t const nChannelsBarrel,
0587 edm::TimeValue_t const event_time,
0588 cudaStream_t cudaStream) {
0589 int nchannels = eventInputGPU.ebUncalibRecHits.size + eventInputGPU.eeUncalibRecHits.size;
0590
0591 unsigned int nchannels_per_block = 16;
0592 unsigned int threads_min = nchannels_per_block;
0593 unsigned int blocks_min = (nchannels + threads_min - 1) / threads_min; // TEST : to be optimized (AM)
0594
0595 //
0596 // kernel create rechit
0597 //
0598
0599 kernel_create_ecal_rehit<<<blocks_min, threads_min, 0, cudaStream>>>(
0600 // configuration
0601 configParameters.ChannelStatusToBeExcluded,
0602 configParameters.ChannelStatusToBeExcludedSize,
0603 configParameters.killDeadChannels,
0604 configParameters.recoverEBIsolatedChannels,
0605 configParameters.recoverEEIsolatedChannels,
0606 configParameters.recoverEBVFE,
0607 configParameters.recoverEEVFE,
0608 configParameters.recoverEBFE,
0609 configParameters.recoverEEFE,
0610 configParameters.EBLaserMIN,
0611 configParameters.EELaserMIN,
0612 configParameters.EBLaserMAX,
0613 configParameters.EELaserMAX,
0614 // for flags setting
0615 configParameters.expanded_v_DB_reco_flags,
0616 configParameters.expanded_Sizes_v_DB_reco_flags,
0617 configParameters.expanded_flagbit_v_DB_reco_flags,
0618 configParameters.expanded_v_DB_reco_flagsSize,
0619 configParameters.flagmask,
0620 // conditions
0621 conditions.ADCToGeV.adc2gev,
0622 conditions.Intercalib.values,
0623 conditions.ChannelStatus.status,
0624 conditions.LaserAPDPNRatiosRef.values,
0625 conditions.LaserAlphas.values,
0626 // input for transparency corrections
0627 conditions.LaserAPDPNRatios.p1,
0628 conditions.LaserAPDPNRatios.p2,
0629 conditions.LaserAPDPNRatios.p3,
0630 conditions.LaserAPDPNRatios.t1,
0631 conditions.LaserAPDPNRatios.t2,
0632 conditions.LaserAPDPNRatios.t3,
0633 // input for linear corrections
0634 conditions.LinearCorrections.p1,
0635 conditions.LinearCorrections.p2,
0636 conditions.LinearCorrections.p3,
0637 conditions.LinearCorrections.t1,
0638 conditions.LinearCorrections.t2,
0639 conditions.LinearCorrections.t3,
0640 // time, used for time dependent corrections
0641 event_time,
0642 // input
0643 eventInputGPU.ebUncalibRecHits.did.get(),
0644 eventInputGPU.eeUncalibRecHits.did.get(),
0645 eventInputGPU.ebUncalibRecHits.amplitude.get(),
0646 eventInputGPU.eeUncalibRecHits.amplitude.get(),
0647 eventInputGPU.ebUncalibRecHits.jitter.get(),
0648 eventInputGPU.eeUncalibRecHits.jitter.get(),
0649 eventInputGPU.ebUncalibRecHits.chi2.get(),
0650 eventInputGPU.eeUncalibRecHits.chi2.get(),
0651 eventInputGPU.ebUncalibRecHits.flags.get(),
0652 eventInputGPU.eeUncalibRecHits.flags.get(),
0653 // output
0654 eventOutputGPU.recHitsEB.did.get(),
0655 eventOutputGPU.recHitsEE.did.get(),
0656 eventOutputGPU.recHitsEB.energy.get(),
0657 eventOutputGPU.recHitsEE.energy.get(),
0658 eventOutputGPU.recHitsEB.time.get(),
0659 eventOutputGPU.recHitsEE.time.get(),
0660 eventOutputGPU.recHitsEB.chi2.get(),
0661 eventOutputGPU.recHitsEE.chi2.get(),
0662 eventOutputGPU.recHitsEB.flagBits.get(),
0663 eventOutputGPU.recHitsEE.flagBits.get(),
0664 eventOutputGPU.recHitsEB.extra.get(),
0665 eventOutputGPU.recHitsEE.extra.get(),
0666 // other
0667 nchannels,
0668 nChannelsBarrel,
0669 conditions.offsetForHashes);
0670 }
0671
0672 } // namespace rechit
0673
0674 } // namespace ecal
