Skip to content

Commit

Permalink
Fix warnings about unused variables in HCAL GPU code (#491)
Browse files Browse the repository at this point in the history
  • Loading branch information
@fwyzard
fwyzard committed Nov 9, 2020
1 parent f907fe1 commit fdc20a6
Show file tree
Hide file tree
Showing 2 changed files with 39 additions and 37 deletions.
38 changes: 17 additions & 21 deletions EventFilter/HcalRawToDigi/plugins/DecodeGPU.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -76,9 +76,7 @@ namespace hcal {

auto const iamc = threadIdx.x / NTHREADS;
auto const ifed = blockIdx.x;
auto const fed = feds[ifed];
auto const offset = offsets[ifed];
auto const size = ifed == gridDim.x - 1 ? nBytesTotal - offset : offsets[ifed + 1] - offset;

#ifdef HCAL_RAWDECODE_GPUDEBUG_CG
if (ifed > 0 || iamc > 0)
Expand All @@ -87,12 +85,15 @@ namespace hcal {
#endif

#ifdef HCAL_RAWDECODE_GPUDEBUG
auto const fed = feds[ifed];
auto const size = ifed == gridDim.x - 1 ? nBytesTotal - offset : offsets[ifed + 1] - offset;
printf("ifed = %d fed = %d offset = %u size = %u\n", ifed, fed, offset, size);
#endif

// offset to the right raw buffer
uint64_t const* buffer = reinterpret_cast<uint64_t const*>(data + offset);

#ifdef HCAL_RAWDECODE_GPUDEBUG
//
// fed header
//
Expand All @@ -103,7 +104,6 @@ namespace hcal {
uint8_t const trigger_type = (fed_header >> 56) & 0xf;
uint8_t const bid_fed_header = (fed_header >> 60) & 0xf;

#ifdef HCAL_RAWDECODE_GPUDEBUG
printf("fed = %d fed_id = %u bx = %u lv1 = %u trigger_type = %u bid = %u\n",
fed,
fed_id,
Expand All @@ -116,13 +116,12 @@ namespace hcal {
// amc 13 header
auto const amc13word = buffer[1];
uint8_t const namc = (amc13word >> 52) & 0xf;
uint8_t const amc13version = (amc13word >> 60) & 0xf;
uint32_t const amc13OrbitNumber = (amc13word >> 4) & 0xffffffffu;

if (iamc >= namc)
return;

#ifdef HCAL_RAWDECODE_GPUDEBUG
uint8_t const amc13version = (amc13word >> 60) & 0xf;
uint32_t const amc13OrbitNumber = (amc13word >> 4) & 0xffffffffu;
printf("fed = %d namc = %u amc13version = %u amc13OrbitNumber = %u\n", fed, namc, amc13version, amc13OrbitNumber);
#endif

Expand All @@ -134,16 +133,14 @@ namespace hcal {
amcoffset += amcSize;
}

// for (uint8_t iamc=0u; iamc < namc; ++iamc) {
auto const word = buffer[2 + iamc];
uint16_t const amcid = word & 0xffff;
int const slot = (word >> 16) & 0xf;
int const amcBlockNumber = (word >> 20) & 0xff;
int const amcSize = (word >> 32) & 0xffffff;

#ifdef HCAL_RAWDECODE_GPUDEBUG
uint16_t const amcid = word & 0xffff;
int const slot = (word >> 16) & 0xf;
int const amcBlockNumber = (word >> 20) & 0xff;
printf("fed = %d amcid = %u slot = %d amcBlockNumber = %d\n", fed, amcid, slot, amcBlockNumber);
#endif

bool const amcmore = ((word >> 61) & 0x1) != 0;
bool const amcSegmented = ((word >> 60) & 0x1) != 0;
Expand All @@ -152,8 +149,6 @@ namespace hcal {
bool const amcDataPresent = ((word >> 58) & 0x1) != 0;
bool const amcDataValid = ((word >> 56) & 0x1) != 0;
bool const amcEnabled = ((word >> 59) & 0x1) != 0;

#ifdef HCAL_RAWDECODE_GPUDEBUG
printf(
"fed = %d amcmore = %d amcSegmented = %d, amcLengthOk = %d amcCROk = %d\n>> amcDataPresent = %d amcDataValid "
"= %d amcEnabled = %d\n",
Expand All @@ -169,18 +164,19 @@ namespace hcal {

// get to the payload
auto const* payload64 = buffer + 2 + namc + amcoffset;
auto const* payload16 = reinterpret_cast<uint16_t const*>(payload64);
//amcoffset += amcSize;

#ifdef HCAL_RAWDECODE_GPUDEBUG
// uhtr header v1 1st 64 bits
auto const payload64_w0 = payload64[0];
// uhtr n bytes comes from amcSize, according to the cpu version!
//uint32_t const data_length64 = payload64_w0 & 0xfffff;
#endif
// uhtr n bytes comes from amcSize, according to the cpu version!
uint32_t const data_length64 = amcSize;
uint16_t bcn = (payload64_w0 >> 20) & 0xfff;
uint32_t evn = (payload64_w0 >> 32) & 0xffffff;

#ifdef HCAL_RAWDECODE_GPUDEBUG
uint16_t bcn = (payload64_w0 >> 20) & 0xfff;
uint32_t evn = (payload64_w0 >> 32) & 0xffffff;
printf("fed = %d data_length64 = %u bcn = %u evn = %u\n", fed, data_length64, bcn, evn);
#endif

Expand All @@ -189,12 +185,12 @@ namespace hcal {
uint8_t const uhtrcrate = payload64_w1 & 0xff;
uint8_t const uhtrslot = (payload64_w1 >> 8) & 0xf;
uint8_t const presamples = (payload64_w1 >> 12) & 0xf;
uint16_t const orbitN = (payload64_w1 >> 16) & 0xffff;
uint8_t const firmFlavor = (payload64_w1 >> 32) & 0xff;
uint8_t const eventType = (payload64_w1 >> 40) & 0xf;
uint8_t const payloadFormat = (payload64_w1 >> 44) & 0xf;

#ifdef HCAL_RAWDECODE_GPUDEBUG
uint16_t const orbitN = (payload64_w1 >> 16) & 0xffff;
uint8_t const firmFlavor = (payload64_w1 >> 32) & 0xff;
uint8_t const eventType = (payload64_w1 >> 40) & 0xf;
printf(
"fed = %d crate = %u slot = %u presamples = %u\n>>> orbitN = %u firmFlavor = %u eventType = %u payloadFormat "
"= %u\n",
Expand All @@ -215,7 +211,7 @@ namespace hcal {
// skip uhtr header words
auto const channelDataSize = data_length64 - 2; // 2 uhtr header v1 words
auto const* channelDataBuffer64Start = payload64 + 2; // 2 uhtr header v2 wds
auto const* channelDataBuffer64End = channelDataBuffer64Start + channelDataSize;
//auto const* channelDataBuffer64End = channelDataBuffer64Start + channelDataSize;
auto const* ptr = reinterpret_cast<uint16_t const*>(channelDataBuffer64Start);
auto const* end = ptr + sizeof(uint64_t) / sizeof(uint16_t) * (channelDataSize - 1);
auto const t_rank = thread_group.thread_rank();
Expand Down
38 changes: 22 additions & 16 deletions RecoLocalCalo/HcalRecProducers/src/MahiGPU.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -28,7 +28,7 @@ namespace hcal {

// FIXME remove duplication...
// this is from PulesFunctor. nvcc was complaining... if included that header...
constexpr int maxSamples = 10;
//constexpr int maxSamples = 10;
constexpr int maxPSshapeBin = 256;
constexpr int nsPerBX = 25;
constexpr float iniTimeShift = 92.5f;
Expand Down Expand Up @@ -229,9 +229,6 @@ namespace hcal {
auto const nchannels_per_block = blockDim.y;
auto const linearThPerBlock = threadIdx.x + threadIdx.y * blockDim.x;

constexpr uint32_t mantissaMaskQIE8 = 0x1fu;
constexpr uint32_t mantissaMaskQIE11 = 0x3f;

// remove
if (gch >= nchannels)
return;
Expand Down Expand Up @@ -375,13 +372,17 @@ namespace hcal {
// sits on the boundary when flavor 01 channels end and flavor 5 start
//
float rawCharge;
#ifdef COMPUTE_TDC_TIME
float tdcTime;
#endif // COMPUTE_TDC_TIME
auto const dfc = compute_diff_charge_gain(
qieType, adc, capid, qieOffsets, qieSlopes, gch < nchannelsf01HE || gch >= nchannelsf015);
if (gch >= nchannelsf01HE && gch < nchannelsf015) {
// flavor 5
rawCharge = charge;
#ifdef COMPUTE_TDC_TIME
tdcTime = HcalSpecialTimes::UNKNOWN_T_NOTDC;
#endif // COMPUTE_TDC_TIME
} else {
// flavor 0 or 1 or 3
// conditions needed for sipms
Expand All @@ -397,11 +398,13 @@ namespace hcal {
auto const effectivePixelsFired = sipmq / fcByPE;
auto const factor = compute_reco_correction_factor(parLin1, parLin2, parLin3, effectivePixelsFired);
rawCharge = (charge - pedestal) * factor + pedestal;
#ifdef COMPUTE_TDC_TIME
if (gch < nchannelsf01HE)
tdcTime = HcalSpecialTimes::getTDCTime(tdc_for_sample<Flavor01>(dataf01HE + stride * gch, sample));
else if (gch >= nchannelsf015)
tdcTime =
HcalSpecialTimes::getTDCTime(tdc_for_sample<Flavor3>(dataf3HB + stride * (gch - nchannelsf015), sample));
#endif // COMPUTE_TDC_TIME

#ifdef HCAL_MAHI_GPUDEBUG
printf("first = %d last = %d sipmQ = %f factor = %f rawCharge = %f\n", first, last, sipmq, factor, rawCharge);
Expand Down Expand Up @@ -454,7 +457,7 @@ namespace hcal {
do {
assumed = old;
// decode energy, sample values
int const current_sample = (assumed >> 32) & 0xffffffff;
//int const current_sample = (assumed >> 32) & 0xffffffff;
float const current_energy = __uint_as_float(assumed & 0xffffffff);
if (energym0_per_ts > current_energy)
old = atomicCAS(&shrMethod0EnergySamplePair[lch], assumed, val);
Expand Down Expand Up @@ -554,8 +557,8 @@ namespace hcal {
constexpr float pulse_height = 1.0f;
constexpr float slew = 0.f;
constexpr auto ns_per_bx = nsPerBX;
constexpr auto num_ns = nsPerBX * maxSamples;
constexpr auto num_bx = num_ns / ns_per_bx;
//constexpr auto num_ns = nsPerBX * maxSamples;
//constexpr auto num_bx = num_ns / ns_per_bx;

// FIXME: clean up all the rounding... this is coming from original cpu version
float const i_start_float =
Expand Down Expand Up @@ -786,9 +789,10 @@ namespace hcal {
// FIXME: shift should be treated properly,
// here assume 8 time slices and 8 samples
auto const shift = 4 - soi; // as in cpu version!
// auto const offset = ipulse - soi;
// auto const idx = sample - offset;
auto const idx = sample - pulseOffset;

// auto const offset = ipulse - soi;
// auto const idx = sample - offset;
int32_t const idx = sample - pulseOffset;
auto const value = idx >= 0 && idx < nsamples ? compute_pulse_shape_value(t0,
idx,
shift,
Expand Down Expand Up @@ -1346,8 +1350,10 @@ namespace hcal {
#endif
#endif

/*
// TODO: provide this properly
int const soi = soiSamples[gch];
*/
constexpr float deltaChi2Threashold = 1e-3;

ColumnVector<NPULSES, int> pulseOffsets;
Expand Down Expand Up @@ -1633,12 +1639,12 @@ namespace hcal {
auto const idx_for_energy = std::abs(pulseOffsetValues[0]);
outputEnergy[gch] = (gain * resultAmplitudesVector(idx_for_energy)) * respCorrection;
/*
#pragma unroll
for (int i=0; i<NPULSES; i++)
if (pulseOffsets[i] == soi)
// NOTE: gain is a number < 10^-3/4, multiply first to avoid stab issues
outputEnergy[gch] = (gain*resultAmplitudesVector(i))*respCorrection;
*/
#pragma unroll
for (int i=0; i<NPULSES; i++)
if (pulseOffsets[i] == soi)
// NOTE: gain is a number < 10^-3/4, multiply first to avoid stab issues
outputEnergy[gch] = (gain*resultAmplitudesVector(i))*respCorrection;
*/
}

void entryPoint(InputDataGPU const& inputGPU,
Expand Down

0 comments on commit fdc20a6

Please sign in to comment.