|
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
|
#ifndef CondFormats_SiPixelObjects_interface_SiPixelGainForHLTonGPU_h
#define CondFormats_SiPixelObjects_interface_SiPixelGainForHLTonGPU_h
#include <cstdint>
#include <cstdio>
#include <tuple>
// including <cuda_runtime.h> would pull in the dependency on all of CUDA;
// instead, just define away the CUDA specific attributes to keep GCC happy.
#ifndef __CUDACC__
#ifndef __host__
#define __host__
#endif // __host__
#ifndef __device__
#define __device__
#endif // __device__
#endif // __CUDACC__
#include "Geometry/CommonTopologies/interface/SimplePixelTopology.h"
#include "HeterogeneousCore/CUDAUtilities/interface/cuda_assert.h"
struct SiPixelGainForHLTonGPU_DecodingStructure {
uint8_t gain;
uint8_t ped;
};
// copy of SiPixelGainCalibrationForHLT
class SiPixelGainForHLTonGPU {
public:
using DecodingStructure = SiPixelGainForHLTonGPU_DecodingStructure;
using Range = std::pair<uint32_t, uint32_t>;
inline __host__ __device__ std::pair<float, float> getPedAndGain(
uint32_t moduleInd, int col, int row, bool& isDeadColumn, bool& isNoisyColumn) const {
auto range = rangeAndCols_[moduleInd].first;
auto nCols = rangeAndCols_[moduleInd].second;
// determine what averaged data block we are in (there should be 1 or 2 of these depending on if plaquette is 1 by X or 2 by X
unsigned int lengthOfColumnData = (range.second - range.first) / nCols;
unsigned int lengthOfAveragedDataInEachColumn = 2; // we always only have two values per column averaged block
unsigned int numberOfDataBlocksToSkip = row / numberOfRowsAveragedOver_;
auto offset = range.first + col * lengthOfColumnData + lengthOfAveragedDataInEachColumn * numberOfDataBlocksToSkip;
assert(offset < range.second);
assert(offset < 3088384);
assert(0 == offset % 2);
DecodingStructure const* __restrict__ lp = v_pedestals_;
auto s = lp[offset / 2];
isDeadColumn = (s.ped & 0xFF) == deadFlag_;
isNoisyColumn = (s.ped & 0xFF) == noisyFlag_;
return std::make_pair(decodePed(s.ped & 0xFF), decodeGain(s.gain & 0xFF));
}
constexpr float decodeGain(unsigned int gain) const { return float(gain) * gainPrecision_ + minGain_; }
constexpr float decodePed(unsigned int ped) const { return float(ped) * pedPrecision_ + minPed_; }
DecodingStructure* v_pedestals_;
std::pair<Range, int> rangeAndCols_[phase1PixelTopology::numberOfModules];
float minPed_, maxPed_, minGain_, maxGain_;
float pedPrecision_, gainPrecision_;
unsigned int numberOfRowsAveragedOver_; // this is 80!!!!
unsigned int nBinsToUseForEncoding_;
unsigned int deadFlag_;
unsigned int noisyFlag_;
};
#endif // CondFormats_SiPixelObjects_interface_SiPixelGainForHLTonGPU_h
|