Skip to content

Commit 3841358

Browse files
fweigdavidrohr
fweig
authored and
davidrohr
committed
GPU/TPCClusterFinder: Add GPU decoder for link based ZS.
1 parent 4806fbf commit 3841358

27 files changed

Lines changed: 592 additions & 46 deletions

‎DataFormats/Detectors/TPC/include/DataFormatsTPC/ZeroSuppression.h‎

Lines changed: 9 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -25,6 +25,12 @@ namespace o2
2525
namespace tpc
2626
{
2727

28+
enum ZSVersion : unsigned char {
29+
ZSVersionRowBased10BitADC = 1,
30+
ZSVersionRowBased12BitADC = 2,
31+
ZSVersionLinkBasedWithMeta = 3,
32+
};
33+
2834
struct TPCZSHDR {
2935
static constexpr size_t TPC_ZS_PAGE_SIZE = 8192;
3036
static constexpr size_t TPC_MAX_SEQ_LEN = 138;
@@ -44,6 +50,9 @@ struct TPCZSHDR {
4450
};
4551
struct TPCZSHDRV2 : public TPCZSHDR {
4652
static constexpr unsigned int TPC_ZS_NBITS_V3 = 12;
53+
static constexpr bool TIGHTLY_PACKED_V3 = false;
54+
static constexpr unsigned int SAMPLESPER64BIT = 64 / TPC_ZS_NBITS_V3; // 5 12-bit samples with 4 bit padding per 64 bit word for non-TIGHTLY_PACKED data
55+
4756
unsigned short firstZSDataOffset; // Offset (after the TPCZSHDRV2 header) in 128bit words to first ZS data (in between can be trigger words, etc.)
4857
unsigned short nTimebinHeaders; // Number of timebin headers
4958
unsigned short reserved1; // 16 reserved bits, header is 128 bit

‎DataFormats/Detectors/TPC/include/DataFormatsTPC/ZeroSuppressionLinkBased.h‎

Lines changed: 11 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -16,7 +16,11 @@
1616
#ifndef ALICEO2_DATAFORMATSTPC_ZeroSuppressionLinkBased_H
1717
#define ALICEO2_DATAFORMATSTPC_ZeroSuppressionLinkBased_H
1818

19+
#include "GPUCommonDef.h"
20+
21+
#ifndef GPUCA_GPUCODE
1922
#include <bitset>
23+
#endif
2024

2125
namespace o2
2226
{
@@ -27,6 +31,7 @@ namespace zerosupp_link_based
2731

2832
static constexpr uint32_t DataWordSizeBits = 128; ///< size of header word and data words in bits
2933
static constexpr uint32_t DataWordSizeBytes = DataWordSizeBits / 8; ///< size of header word and data words in bytes
34+
static constexpr uint32_t ChannelPerTBHeader = 80;
3035

3136
/// common header definition of the zero suppressed link based data
3237
struct CommonHeader {
@@ -61,6 +66,9 @@ struct CommonHeader {
6166
bool isTriggerInfoV2() const { return (magicWord == MagicWordTriggerV2); }
6267
};
6368

69+
// GPU Code only requires header definition
70+
#ifndef GPUCA_GPUCODE
71+
6472
/// header definition of the zero suppressed link based data format
6573
struct Header final : public CommonHeader {
6674

@@ -273,6 +281,9 @@ struct TriggerContainer {
273281

274282
uint32_t getTriggerType() const { return triggerInfo.getTriggerType(); }
275283
};
284+
285+
#endif // !defined(GPUCA_GPUCODE)
286+
276287
} // namespace zerosupp_link_based
277288
} // namespace tpc
278289
} // namespace o2

‎Detectors/TPC/base/include/TPCBase/Mapper.h‎

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -436,6 +436,9 @@ class Mapper
436436
return getPadsInOROC();
437437
}
438438

439+
const std::vector<PadPos>& getMapGlobalPadToPadPos() const { return mMapGlobalPadToPadPos; }
440+
const std::vector<int>& getMapFECIDGlobalPad() const { return mMapFECIDGlobalPad; }
441+
439442
const std::vector<float>& getTraceLengthsIROC() const { return mTraceLengthsIROC; }
440443
const std::vector<float>& getTraceLengthsOROC() const { return mTraceLengthsOROC; }
441444

‎Detectors/TPC/base/include/TPCBase/PadPos.h‎

Lines changed: 13 additions & 11 deletions
Original file line numberDiff line numberDiff line change
@@ -25,6 +25,8 @@
2525
#ifndef AliceO2_TPC_PadPos_H
2626
#define AliceO2_TPC_PadPos_H
2727

28+
#include "GPUCommonDef.h"
29+
2830
namespace o2
2931
{
3032
namespace tpc
@@ -38,49 +40,49 @@ class PadPos
3840
/// constructor
3941
/// @param [in] row pad row
4042
/// @param [in] pad pad in a row
41-
PadPos(const unsigned char row, const unsigned char pad) : mRow(row), mPad(pad) {}
43+
GPUdi() PadPos(const unsigned char row, const unsigned char pad) : mRow(row), mPad(pad) {}
4244

4345
/// numeric row number
4446
/// @return numeric row number
45-
unsigned char getRow() const { return mRow; }
47+
GPUdi() unsigned char getRow() const { return mRow; }
4648

4749
/// numeric pad number
4850
/// @return numeric pad number
49-
unsigned char getPad() const { return mPad; }
51+
GPUdi() unsigned char getPad() const { return mPad; }
5052

5153
/// setter for row number
5254
/// @param [in] row row number
53-
void setRow(const unsigned char row) { mRow = row; }
55+
GPUdi() void setRow(const unsigned char row) { mRow = row; }
5456

5557
/// setter for pad number
5658
/// @param [in] pad pad number
57-
void setPad(const unsigned char pad) { mPad = pad; }
59+
GPUdi() void setPad(const unsigned char pad) { mPad = pad; }
5860

5961
/// add row offset
6062
/// @param [in] rowOffset row offset to add
61-
void addRowOffset(const unsigned char rowOffset) { mRow += rowOffset; }
63+
GPUdi() void addRowOffset(const unsigned char rowOffset) { mRow += rowOffset; }
6264

6365
/// setter for row and pad number
6466
/// @param [in] row row number
6567
/// @param [in] pad pad number
66-
void set(const unsigned char row, const unsigned char pad)
68+
GPUdi() void set(const unsigned char row, const unsigned char pad)
6769
{
6870
mRow = row;
6971
mPad = pad;
7072
}
7173

7274
/// check if is valid
7375
/// @return pad valid
74-
bool isValid() const { return !(mRow == 255 && mPad == 255); }
76+
GPUdi() bool isValid() const { return !(mRow == 255 && mPad == 255); }
7577

7678
/// equal operator
77-
bool operator==(const PadPos& other) const { return (mRow == other.mRow) && (mPad == other.mPad); }
79+
GPUdi() bool operator==(const PadPos& other) const { return (mRow == other.mRow) && (mPad == other.mPad); }
7880

7981
/// unequal operator
80-
bool operator!=(const PadPos& other) const { return (mRow != other.mRow) || (mPad != other.mPad); }
82+
GPUdi() bool operator!=(const PadPos& other) const { return (mRow != other.mRow) || (mPad != other.mPad); }
8183

8284
/// smaller operator
83-
bool operator<(const PadPos& other) const
85+
GPUdi() bool operator<(const PadPos& other) const
8486
{
8587
if (mRow < other.mRow) {
8688
return true;

‎GPU/Common/GPUCommonAlgorithm.h‎

Lines changed: 21 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -340,6 +340,10 @@ GPUdi() void GPUCommonAlgorithm::swap(T& a, T& b)
340340

341341
#ifdef __OPENCL__
342342
// Nothing to do, work_group functions available
343+
#pragma OPENCL EXTENSION cl_khr_subgroups : enable
344+
345+
#define warp_scan_inclusive_add(v) sub_group_scan_inclusive_add(v)
346+
#define warp_broadcast(v, i) sub_group_broadcast(v, i)
343347

344348
#elif (defined(__CUDACC__) || defined(__HIPCC__))
345349
// CUDA and HIP work the same way using cub, need just different header
@@ -392,6 +396,17 @@ GPUdi() T warp_scan_inclusive_add_FUNC(T v, S& smem)
392396
return v;
393397
}
394398

399+
#define warp_broadcast(v, i) warp_broadcast_FUNC(v, i, smem)
400+
template <class T, class S>
401+
GPUdi() T warp_broadcast_FUNC(T v, int i, S& smem)
402+
{
403+
#ifdef __CUDACC__
404+
return __shfl_sync(0xFFFFFFFF, v, i);
405+
#else // HIP
406+
return __shfl(v, i);
407+
#endif
408+
}
409+
395410
#else
396411
// Trivial implementation for the CPU
397412

@@ -419,6 +434,12 @@ GPUdi() T warp_scan_inclusive_add(T v)
419434
return v;
420435
}
421436

437+
template <class T>
438+
GPUdi() T warp_broadcast(T v, int i)
439+
{
440+
return v;
441+
}
442+
422443
#endif
423444

424445
#endif

‎GPU/GPUTracking/Base/GPUGeneralKernels.h‎

Lines changed: 11 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -54,6 +54,17 @@ class GPUKernelTemplate
5454
struct GPUSharedMemory {
5555
};
5656

57+
template <class T, int I>
58+
struct GPUSharedMemoryWarpScan64 {
59+
// Provides the shared memory resources for warp wide CUB collectives
60+
#if (defined(__CUDACC__) || defined(__HIPCC__)) && defined(GPUCA_GPUCODE) && !defined(GPUCA_GPUCODE_HOSTONLY)
61+
typedef GPUCA_CUB::WarpScan<T> WarpScan;
62+
union {
63+
typename WarpScan::TempStorage cubWarpTmpMem;
64+
};
65+
#endif
66+
};
67+
5768
template <class T, int I>
5869
struct GPUSharedMemoryScan64 {
5970
// Provides the shared memory resources for CUB collectives

‎GPU/GPUTracking/Base/GPUReconstructionConvert.cxx‎

Lines changed: 7 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -457,9 +457,6 @@ void zsEncoderRow::decodePage(std::vector<o2::tpc::Digit>& outputBuffer, const z
457457

458458
#ifdef GPUCA_O2_LIB
459459
struct zsEncoderImprovedLinkBased : public zsEncoder {
460-
static constexpr bool TIGHTLY_PACKED = false;
461-
static constexpr unsigned int SAMPLESPER64BIT = 64 / TPCZSHDRV2::TPC_ZS_NBITS_V3; // 5 12-bit samples with 4 bit padding per 64 bit word for non-TIGHTLY_PACKED data
462-
463460
TPCZSHDRV2* hdr = nullptr;
464461
int inverseChannelMapping[5][32];
465462
int nSamples = 0;
@@ -545,10 +542,10 @@ bool zsEncoderImprovedLinkBased::checkInput(std::vector<o2::tpc::Digit>& tmpBuff
545542
if (!finishPage) {
546543
unsigned int sizeChk = (unsigned int)(pagePtr - reinterpret_cast<unsigned char*>(page));
547544
sizeChk += sizeof(o2::tpc::zerosupp_link_based::CommonHeader);
548-
if (TIGHTLY_PACKED) {
545+
if (TPCZSHDRV2::TIGHTLY_PACKED_V3) {
549546
sizeChk += (nSamples * TPCZSHDRV2::TPC_ZS_NBITS_V3 + 127) / 128 * 16;
550547
} else {
551-
sizeChk += (nSamples + 2 * SAMPLESPER64BIT - 1) / (2 * SAMPLESPER64BIT) * 16;
548+
sizeChk += (nSamples + 2 * TPCZSHDRV2::SAMPLESPER64BIT - 1) / (2 * TPCZSHDRV2::SAMPLESPER64BIT) * 16;
552549
}
553550
if (sizeChk > TPCZSHDR::TPC_ZS_PAGE_SIZE) {
554551
finishPage = true;
@@ -574,19 +571,19 @@ unsigned int zsEncoderImprovedLinkBased::encodeSequence(std::vector<o2::tpc::Dig
574571
tbHdr->fecInPartition = link;
575572
hdr->nTimeBins = tmpBuffer[k].getTimeStamp() - firstTimebinInPage + 1;
576573
hdr->nTimebinHeaders++;
577-
if (TIGHTLY_PACKED) {
574+
if (TPCZSHDRV2::TIGHTLY_PACKED_V3) {
578575
tbHdr->numWordsPayload = (nSamples * TPCZSHDRV2::TPC_ZS_NBITS_V3 + 127) / 128; // tightly packed ADC samples
579576
unsigned int tmp = 0;
580577
unsigned int tmpIn = nSamples;
581578
ZSstreamOut(adcValues.data(), tmpIn, pagePtr, tmp, encodeBits);
582579
} else {
583-
tbHdr->numWordsPayload = (nSamples + 2 * SAMPLESPER64BIT - 1) / (2 * SAMPLESPER64BIT);
580+
tbHdr->numWordsPayload = (nSamples + 2 * TPCZSHDRV2::SAMPLESPER64BIT - 1) / (2 * TPCZSHDRV2::SAMPLESPER64BIT);
584581
unsigned long* payloadPtr = (unsigned long*)pagePtr;
585582
for (unsigned int i = 0; i < 2 * tbHdr->numWordsPayload; i++) {
586583
payloadPtr[i] = 0;
587584
}
588585
for (unsigned int i = 0; i < nSamples; i++) {
589-
payloadPtr[i / SAMPLESPER64BIT] |= ((unsigned long)adcValues[i]) << ((i % SAMPLESPER64BIT) * TPCZSHDRV2::TPC_ZS_NBITS_V3);
586+
payloadPtr[i / TPCZSHDRV2::SAMPLESPER64BIT] |= ((unsigned long)adcValues[i]) << ((i % TPCZSHDRV2::SAMPLESPER64BIT) * TPCZSHDRV2::TPC_ZS_NBITS_V3);
590587
}
591588
}
592589
pagePtr += tbHdr->numWordsPayload * 16;
@@ -670,7 +667,7 @@ void zsEncoderImprovedLinkBased::decodePage(std::vector<o2::tpc::Digit>& outputB
670667
const auto& bitmask = tbHdr->getChannelBits();
671668
int nADC = bitmask.count();
672669
std::vector<unsigned short> decBuffer(nADC);
673-
if (TIGHTLY_PACKED) {
670+
if (TPCZSHDRV2::TIGHTLY_PACKED_V3) {
674671
unsigned int byte = 0, bits = 0, posXbits = 0;
675672
while (posXbits < nADC) {
676673
byte |= *(adcData++) << bits;
@@ -684,7 +681,7 @@ void zsEncoderImprovedLinkBased::decodePage(std::vector<o2::tpc::Digit>& outputB
684681
} else {
685682
const unsigned long* adcData64 = (const unsigned long*)adcData;
686683
for (int j = 0; j < nADC; j++) {
687-
decBuffer[j] = (adcData64[j / SAMPLESPER64BIT] >> ((j % SAMPLESPER64BIT) * TPCZSHDRV2::TPC_ZS_NBITS_V3)) & mask;
684+
decBuffer[j] = (adcData64[j / TPCZSHDRV2::SAMPLESPER64BIT] >> ((j % TPCZSHDRV2::SAMPLESPER64BIT) * TPCZSHDRV2::TPC_ZS_NBITS_V3)) & mask;
688685
}
689686
}
690687
for (int j = 0, k = 0; j < bitmask.size(); j++) {

‎GPU/GPUTracking/Base/GPUReconstructionKernels.h‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -94,6 +94,7 @@ GPUCA_KRNL_LB((GPUTPCCFStreamCompaction, scanTop ), (single), (, int
9494
GPUCA_KRNL_LB((GPUTPCCFStreamCompaction, scanDown ), (single), (, int iBuf, unsigned int offset, int nElems), (, iBuf, offset, nElems))
9595
GPUCA_KRNL_LB((GPUTPCCFStreamCompaction, compactDigits ), (single), (, int iBuf, int stage, GPUPtr1(ChargePos*, in), GPUPtr1(ChargePos*, out)), (, iBuf, stage, GPUPtr2(ChargePos*, in), GPUPtr2(ChargePos*, out)))
9696
GPUCA_KRNL_LB((GPUTPCCFDecodeZS ), (single), (, int firstHBF), (, firstHBF))
97+
GPUCA_KRNL_LB((GPUTPCCFDecodeZSLink ), (single), (, int firstHBF), (, firstHBF))
9798
GPUCA_KRNL_LB((GPUTPCCFGather ), (single), (, GPUPtr1(o2::tpc::ClusterNative*, dest)), (, GPUPtr2(o2::tpc::ClusterNative*, dest)))
9899
GPUCA_KRNL_LB((GPUTrackingRefitKernel, mode0asGPU ), (simple), (), ())
99100
GPUCA_KRNL_LB((GPUTrackingRefitKernel, mode1asTrackParCov ), (simple), (), ())

‎GPU/GPUTracking/CMakeLists.txt‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -205,6 +205,7 @@ if(ALIGPU_BUILD_TYPE STREQUAL "O2" OR CONFIG_O2_EXTENSIONS)
205205
TPCClusterFinder/GPUTPCCFDecodeZS.cxx
206206
TPCClusterFinder/GPUTPCCFGather.cxx
207207
DataTypes/TPCPadGainCalib.cxx
208+
DataTypes/TPCZSLinkMapping.cxx
208209
DataTypes/CalibdEdxContainer.cxx
209210
Refit/GPUTrackingRefit.cxx
210211
Refit/GPUTrackingRefitKernel.cxx

‎GPU/GPUTracking/DataTypes/GPUDataTypes.h‎

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -105,7 +105,7 @@ namespace gpu
105105
{
106106
class TPCFastTransform;
107107
struct TPCPadGainCalib;
108-
108+
struct TPCZSLinkMapping;
109109
} // namespace gpu
110110
} // namespace GPUCA_NAMESPACE
111111

@@ -215,6 +215,7 @@ struct GPUCalibObjectsTemplate {
215215
typename S<o2::base::MatLayerCylSet>::type* matLUT = nullptr;
216216
typename S<o2::trd::GeometryFlat>::type* trdGeometry = nullptr;
217217
typename S<TPCPadGainCalib>::type* tpcPadGain = nullptr;
218+
typename S<TPCZSLinkMapping>::type* tpcZSLinkMapping = nullptr;
218219
typename S<o2::tpc::CalibdEdxContainer>::type* dEdxCalibContainer = nullptr;
219220
typename S<o2::base::PropagatorImpl<float>>::type* o2Propagator = nullptr;
220221
typename S<o2::itsmft::TopologyDictionary>::type* itsPatternDict = nullptr;

0 commit comments

Comments
 (0)