Project
Loading...
Searching...
No Matches
GPUTPCCFDeconvolution.cxx
Go to the documentation of this file.
1// Copyright 2019-2020 CERN and copyright holders of ALICE O2.
2// See https://alice-o2.web.cern.ch/copyright for details of the copyright holders.
3// All rights not expressly granted are reserved.
4//
5// This software is distributed under the terms of the GNU General Public
6// License v3 (GPL Version 3), copied verbatim in the file "COPYING".
7//
8// In applying this license CERN does not waive the privileges and immunities
9// granted to it by virtue of its status as an Intergovernmental Organization
10// or submit itself to any jurisdiction.
11
14
16#include "CfConsts.h"
17#include "CfUtils.h"
18#include "CfChargePos.h"
19#include "GPUDefMacros.h"
20
21using namespace o2::gpu;
22using namespace o2::gpu::tpccf;
23
24template <>
25GPUdii() void GPUTPCCFDeconvolution::Thread<0>(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUSharedMemory& smem, processorType& clusterer, uint8_t overwriteCharge)
26{
27 CfArray2D<PackedCharge> chargeMap(reinterpret_cast<PackedCharge*>(clusterer.mPchargeMap));
28 CfArray2D<uint8_t> isPeakMap(clusterer.mPpeakMap);
29 GPUTPCCFDeconvolution::deconvolutionImpl(get_num_groups(0), get_local_size(0), get_group_id(0), get_local_id(0), smem, isPeakMap, chargeMap, clusterer.mPpositions, clusterer.mPmemory->counters.nPositions, overwriteCharge);
30}
31
32GPUdii() void GPUTPCCFDeconvolution::deconvolutionImpl(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUSharedMemory& smem,
33 const CfArray2D<uint8_t>& peakMap,
34 CfArray2D<PackedCharge>& chargeMap,
35 const CfChargePos* positions,
36 const uint32_t digitnum,
37 uint8_t overwriteCharge)
38{
40
41 bool iamDummy = (idx >= digitnum);
42 idx = iamDummy ? digitnum - 1 : idx;
43
44 CfChargePos pos = positions[idx];
45
46 bool iamPeak = CfUtils::isPeak(peakMap[pos]);
47
48 int8_t peakCount = (iamPeak) ? 1 : 0;
49
50 uint16_t ll = get_local_id(0);
51 uint16_t partId = ll;
52
53 uint16_t in3x3 = 0;
54 bool exclude3x3 = iamPeak || !pos.valid();
55 partId = CfUtils::partition<SCRATCH_PAD_WORK_GROUP_SIZE>(smem, ll, exclude3x3, SCRATCH_PAD_WORK_GROUP_SIZE, &in3x3);
56
57 if (partId < in3x3) {
58 smem.posBcast1[partId] = pos;
59 }
60 GPUbarrier();
61
62 CfUtils::blockLoad(
63 peakMap,
64 in3x3,
65 SCRATCH_PAD_WORK_GROUP_SIZE,
66 ll,
67 0,
68 8,
69 cfconsts::InnerNeighbors,
70 smem.posBcast1,
71 smem.buf);
72
73 uint8_t aboveThreshold = 0;
74 if (partId < in3x3) {
75 peakCount = countPeaksInner(partId, smem.buf, &aboveThreshold);
76 }
77
78 uint16_t in5x5 = 0;
79 partId = CfUtils::partition<SCRATCH_PAD_WORK_GROUP_SIZE>(smem, partId, peakCount > 0 && !exclude3x3, in3x3, &in5x5);
80
81 if (partId < in5x5) {
82 smem.posBcast1[partId] = pos;
83 smem.aboveThresholdBcast[partId] = aboveThreshold;
84 }
85 GPUbarrier();
86
87 CfUtils::condBlockLoad<uint8_t, true>(
88 peakMap,
89 in5x5,
90 SCRATCH_PAD_WORK_GROUP_SIZE,
91 ll,
92 0,
93 16,
94 cfconsts::OuterNeighbors,
95 smem.posBcast1,
96 smem.aboveThresholdBcast,
97 smem.buf);
98
99 if (partId < in5x5) {
100 peakCount = countPeaksOuter(partId, aboveThreshold, smem.buf);
101 peakCount *= -1;
102 }
103
104 if (iamDummy || !pos.valid()) {
105 return;
106 }
107
108 bool has3x3 = (peakCount > 0);
109 peakCount = CAMath::Abs(int32_t(peakCount));
110 bool split = (peakCount > 1);
111
112 peakCount = (peakCount == 0) ? 1 : peakCount;
113
114 PackedCharge charge = chargeMap[pos];
115
116 if (overwriteCharge) {
117 PackedCharge p(charge.unpack() / peakCount, has3x3, split);
118 chargeMap[pos] = p;
119 } else {
120 PackedCharge p(charge.unpack(), has3x3, split);
121 chargeMap[pos] = p;
122 }
123}
124
125GPUdi() uint8_t GPUTPCCFDeconvolution::countPeaksInner(
126 uint16_t ll,
127 const uint8_t* isPeak,
128 uint8_t* aboveThreshold)
129{
130 uint8_t peaks = 0;
131 GPUCA_UNROLL(U(), U())
132 for (uint8_t i = 0; i < 8; i++) {
133 uint8_t p = isPeak[ll * 8 + i];
134 peaks += CfUtils::isPeak(p);
135 *aboveThreshold |= uint8_t(CfUtils::isAboveThreshold(p)) << i;
136 }
137
138 return peaks;
139}
140
141GPUdi() uint8_t GPUTPCCFDeconvolution::countPeaksOuter(
142 uint16_t ll,
143 uint8_t aboveThreshold,
144 const uint8_t* isPeak)
145{
146 uint8_t peaks = 0;
147 GPUCA_UNROLL(U(), U())
148 for (uint8_t i = 0; i < 16; i++) {
149 uint8_t p = isPeak[ll * 16 + i];
150 peaks += CfUtils::isPeak(p);
151 }
152
153 return peaks;
154}
int16_t charge
Definition RawEventData.h:5
int32_t i
#define get_local_size(dim)
#define get_local_id(dim)
#define get_num_groups(dim)
#define GPUbarrier()
#define get_global_id(dim)
#define get_group_id(dim)
#define GPUCA_UNROLL(optCu, optHi)
GPUdii() void GPUTPCCFDeconvolution
uint16_t pos
Definition RawData.h:3
typedef void(APIENTRYP PFNGLCULLFACEPROC)(GLenum mode)
uint8_t itsSharedClusterMap uint8_t
GPUdi() o2
Definition TrackTRD.h:38
std::vector< std::string > split(const std::string &str, char delimiter=',')
for(int irof=0;irof< 1000;irof++)