22#include "GPUTPCCompressionKernels.inc"
42 bool rejectTrk = CAMath::Abs(trk.GetParam().GetQPt() * processors.param.qptB5Scaler) > processors.param.rec.tpc.rejectQPtB5 || trk.MergedLooper();
43 uint32_t nClustersStored = 0;
45 uint8_t lastRow = 0, lastSector = 0;
48 for (int32_t k = trk.NClusters() - 1; k >= 0; k--) {
54 int32_t hitId = hit.
num;
70 float x = geo.Row2X(hit.
row);
71 float y = track.LinearPad2Y(hit.
sector, orgCl.getPad(), geo.PadWidth(hit.
row), geo.NPads(hit.
row));
72 float z = geo.LinearTime2Z(hit.
sector, orgCl.getTime());
73 if (nClustersStored) {
77 if (track.Propagate(geo.Row2X(hit.
row),
param.SectorParam[hit.
sector].Alpha)) {
84 int32_t cidx = trk.FirstClusterRef() + nClustersStored++;
85 if (nClustersStored == 1) {
86 uint8_t qpt = fabs(trk.GetParam().GetQPt()) < 20.f ? (trk.GetParam().GetQPt() * (127.f / 20.f) + 127.5f) : (trk.GetParam().GetQPt() > 0 ? 254 : 0);
92 c.qPtA[myTrack] = qpt;
93 c.rowA[myTrack] = hit.
row;
95 c.timeA[myTrack] = orgCl.getTimePacked();
99 uint32_t sector = hit.
sector;
106 if (lastSector > sector) {
109 sector -= lastSector;
111 c.rowDiffA[cidx] =
row;
112 c.sliceLegDiffA[cidx] = sector;
113 float pad = CAMath::Max(0.f, CAMath::Min((
float)geo.NPads(
GPUCA_ROW_COUNT - 1), track.LinearY2Pad(hit.
sector, track.Y(), geo.PadWidth(hit.
row), geo.NPads(hit.
row))));
114 c.padResA[cidx] = orgCl.
padPacked - orgCl.packPad(pad);
115 float time = CAMath::Max(0.f, geo.LinearZ2Time(hit.
sector, track.Z() + zOffset));
116 c.timeResA[cidx] = (orgCl.getTimePacked() - orgCl.packTime(
time)) & 0xFFFFFF;
118 uint16_t qtot = orgCl.
qTot, qmax = orgCl.
qMax;
121 compressor.truncateSignificantBitsChargeMax(qmax,
param);
122 compressor.truncateSignificantBitsCharge(qtot,
param);
123 compressor.truncateSignificantBitsWidth(sigmapad,
param);
124 compressor.truncateSignificantBitsWidth(sigmatime,
param);
126 c.qTotA[cidx] = qtot;
127 c.qMaxA[cidx] = qmax;
128 c.sigmaPadA[cidx] = sigmapad;
129 c.sigmaTimeA[cidx] = sigmatime;
130 c.flagsA[cidx] = orgCl.getFlags();
131 if (k && track.Filter(
y,
z - zOffset, hit.
row)) {
137 if (nClustersStored) {
139 c.nTrackClusters[myTrack] = nClustersStored;
147 return mClsPtr[
a].getTimePacked() < mClsPtr[
b].getTimePacked();
153 return mClsPtr[
a].padPacked < mClsPtr[
b].padPacked;
159 if (mClsPtr[
a].getTimePacked() >> 3 == mClsPtr[
b].getTimePacked() >> 3) {
160 return mClsPtr[
a].padPacked < mClsPtr[
b].padPacked;
162 return mClsPtr[
a].getTimePacked() < mClsPtr[
b].getTimePacked();
168 if (mClsPtr[
a].padPacked >> 3 == mClsPtr[
b].padPacked >> 3) {
169 return mClsPtr[
a].getTimePacked() < mClsPtr[
b].getTimePacked();
171 return mClsPtr[
a].padPacked < mClsPtr[
b].padPacked;
177 if (mClsPtr[
a].getTimePacked() != mClsPtr[
b].getTimePacked()) {
178 return mClsPtr[
a].getTimePacked() < mClsPtr[
b].getTimePacked();
180 if (mClsPtr[
a].padPacked != mClsPtr[
b].padPacked) {
181 return mClsPtr[
a].padPacked < mClsPtr[
b].padPacked;
183 return mClsPtr[
a].qTot < mClsPtr[
b].qTot;
193 uint32_t* sortBuffer = smem.sortBuffer;
197 const uint32_t idOffset =
clusters->clusterOffset[iSector][iRow];
198 const uint32_t idOffsetOut =
clusters->clusterOffset[iSector][iRow] * compressor.mMaxClusterFactorBase1024 / 1024;
199 const uint32_t idOffsetOutMax = ((
const uint32_t*)
clusters->clusterOffset[iSector])[iRow + 1] * compressor.mMaxClusterFactorBase1024 / 1024;
200 if (iThread == nThreads - 1) {
203 uint32_t totalCount = 0;
209 for (uint32_t
i = iThread;
i < nn + nThreads;
i += nThreads) {
210 const int32_t
idx = idOffset +
i;
211 int32_t storeCluster = 0;
213 if (
i >=
clusters->nClusters[iSector][iRow]) {
216 if (compressor.mClusterStatus[idx]) {
219 int32_t attach = ioPtrs.mergedTrackHitAttachment[
idx];
220 bool unattached = attach == 0;
231 auto& trk = ioPtrs.mergedTracks[
id];
232 if (CAMath::Abs(trk.GetParam().GetQPt() * processors.param.qptB5Scaler) > processors.param.rec.tpc.rejectQPtB5 || trk.MergedLooper()) {
240 int32_t myIndex = work_group_scan_inclusive_add(storeCluster);
241 int32_t storeLater = -1;
244 sortBuffer[smem.nCount + myIndex - 1] =
i;
250 if (iThread == nThreads - 1) {
251 smem.nCount += myIndex;
260 if (idOffsetOut + totalCount +
count > idOffsetOutMax) {
261 if (iThread == nThreads - 1) {
262 compressor.raiseError(GPUErrors::ERROR_COMPRESSION_ROW_HIT_OVERFLOW, iSector * 1000 + iRow, idOffsetOut + totalCount +
count, idOffsetOutMax);
270#ifdef GPUCA_DETERMINISTIC_MODE
271 CAAlgo::sortInBlock(sortBuffer, sortBuffer +
count, GPUTPCCompressionKernels_Compare<GPUSettings::SortZPadTime>(
clusters->clusters[iSector][iRow]));
274 CAAlgo::sortInBlock(sortBuffer, sortBuffer +
count, GPUTPCCompressionKernels_Compare<GPUSettings::SortZPadTime>(
clusters->clusters[iSector][iRow]));
276 CAAlgo::sortInBlock(sortBuffer, sortBuffer +
count, GPUTPCCompressionKernels_Compare<GPUSettings::SortZTimePad>(
clusters->clusters[iSector][iRow]));
278 CAAlgo::sortInBlock(sortBuffer, sortBuffer +
count, GPUTPCCompressionKernels_Compare<GPUSettings::SortPad>(
clusters->clusters[iSector][iRow]));
280 CAAlgo::sortInBlock(sortBuffer, sortBuffer +
count, GPUTPCCompressionKernels_Compare<GPUSettings::SortTime>(
clusters->clusters[iSector][iRow]));
287 int32_t outidx = idOffsetOut + totalCount +
j;
290 int32_t preId =
j != 0 ? (int32_t)sortBuffer[
j - 1] : (totalCount != 0 ? (int32_t)smem.lastIndex : -1);
291 GPUTPCCompression_EncodeUnattached(
param.
rec.tpc.compressionTypeMask, orgCl,
c.timeDiffU[outidx],
c.padDiffU[outidx], preId == -1 ?
nullptr : &
clusters->
clusters[iSector][iRow][preId]);
293 uint16_t qtot = orgCl.qTot, qmax = orgCl.qMax;
294 uint8_t sigmapad = orgCl.sigmaPadPacked, sigmatime = orgCl.sigmaTimePacked;
296 compressor.truncateSignificantBitsChargeMax(qmax,
param);
297 compressor.truncateSignificantBitsCharge(qtot,
param);
298 compressor.truncateSignificantBitsWidth(sigmapad,
param);
299 compressor.truncateSignificantBitsWidth(sigmatime,
param);
301 c.qTotU[outidx] = qtot;
302 c.qMaxU[outidx] = qmax;
303 c.sigmaPadU[outidx] = sigmapad;
304 c.sigmaTimeU[outidx] = sigmatime;
305 c.flagsU[outidx] = orgCl.getFlags();
309 if (storeLater >= 0) {
310 sortBuffer[storeLater] =
i;
313 if (iThread == nThreads - 1 &&
count) {
314 smem.lastIndex = sortBuffer[
count - 1];
315 smem.nCount -=
count;
319 if (iThread == nThreads - 1) {
321 CAMath::AtomicAdd(&compressor.mMemory->nStoredUnattachedClusters, totalCount);
342 return buf128[iWarp];
345template <
typename T,
typename S>
348 if constexpr (
alignof(
S) >=
alignof(T)) {
349 static_cast<void>(
ptr);
352 return reinterpret_cast<size_t>(
ptr) %
alignof(T) == 0;
359 constexpr const int32_t vec128Elems = CpyVector<uint8_t, Vec128>::Size;
360 constexpr const int32_t vec64Elems = CpyVector<uint8_t, Vec64>::Size;
361 constexpr const int32_t vec32Elems = CpyVector<uint8_t, Vec32>::Size;
362 constexpr const int32_t vec16Elems = CpyVector<uint8_t, Vec16>::Size;
364 if (
size >= uint32_t(nThreads * vec128Elems)) {
365 compressorMemcpyVectorised<uint8_t, Vec128>(
dst,
src,
size, nThreads, iThread);
366 }
else if (
size >= uint32_t(nThreads * vec64Elems)) {
367 compressorMemcpyVectorised<uint8_t, Vec64>(
dst,
src,
size, nThreads, iThread);
368 }
else if (
size >= uint32_t(nThreads * vec32Elems)) {
369 compressorMemcpyVectorised<uint8_t, Vec32>(
dst,
src,
size, nThreads, iThread);
370 }
else if (
size >= uint32_t(nThreads * vec16Elems)) {
371 compressorMemcpyVectorised<uint8_t, Vec16>(
dst,
src,
size, nThreads, iThread);
373 compressorMemcpyBasic(
dst,
src,
size, nThreads, iThread);
380 constexpr const int32_t vec128Elems = CpyVector<uint16_t, Vec128>::Size;
381 constexpr const int32_t vec64Elems = CpyVector<uint16_t, Vec64>::Size;
382 constexpr const int32_t vec32Elems = CpyVector<uint16_t, Vec32>::Size;
384 if (
size >= uint32_t(nThreads * vec128Elems)) {
385 compressorMemcpyVectorised<uint16_t, Vec128>(
dst,
src,
size, nThreads, iThread);
386 }
else if (
size >= uint32_t(nThreads * vec64Elems)) {
387 compressorMemcpyVectorised<uint16_t, Vec64>(
dst,
src,
size, nThreads, iThread);
388 }
else if (
size >= uint32_t(nThreads * vec32Elems)) {
389 compressorMemcpyVectorised<uint16_t, Vec32>(
dst,
src,
size, nThreads, iThread);
391 compressorMemcpyBasic(
dst,
src,
size, nThreads, iThread);
398 constexpr const int32_t vec128Elems = CpyVector<uint32_t, Vec128>::Size;
399 constexpr const int32_t vec64Elems = CpyVector<uint32_t, Vec64>::Size;
401 if (
size >= uint32_t(nThreads * vec128Elems)) {
402 compressorMemcpyVectorised<uint32_t, Vec128>(
dst,
src,
size, nThreads, iThread);
403 }
else if (
size >= uint32_t(nThreads * vec64Elems)) {
404 compressorMemcpyVectorised<uint32_t, Vec64>(
dst,
src,
size, nThreads, iThread);
406 compressorMemcpyBasic(
dst,
src,
size, nThreads, iThread);
410template <
typename Scalar,
typename BaseVector>
413 if (not isAlignedTo<BaseVector>(
dst)) {
414 size_t dsti =
reinterpret_cast<size_t>(
dst);
415 int32_t
offset = (
alignof(BaseVector) - dsti %
alignof(BaseVector)) /
sizeof(Scalar);
416 compressorMemcpyBasic(
dst,
src,
offset, nThreads, iThread);
422 BaseVector*
GPUrestrict() dstAligned = reinterpret_cast<BaseVector*>(
dst);
424 using CpyVec = CpyVector<Scalar, BaseVector>;
425 uint32_t sizeAligned =
size / CpyVec::
Size;
427 if (isAlignedTo<BaseVector>(
src)) {
428 const BaseVector*
GPUrestrict() srcAligned = reinterpret_cast<const BaseVector*>(
src);
429 compressorMemcpyBasic(dstAligned, srcAligned, sizeAligned, nThreads, iThread);
431 for (uint32_t
i = iThread;
i < sizeAligned;
i += nThreads) {
433 for (uint32_t
j = 0;
j < CpyVec::Size;
j++) {
434 buf.elems[
j] =
src[
i * CpyVec::Size +
j];
436 dstAligned[
i] =
buf.all;
440 int32_t leftovers =
size % CpyVec::Size;
441 compressorMemcpyBasic(
dst +
size - leftovers,
src +
size - leftovers, leftovers, nThreads, iThread);
447 uint32_t
start = (
size + nBlocks - 1) / nBlocks * iBlock + iThread;
448 uint32_t
end = CAMath::Min(
size, (
size + nBlocks - 1) / nBlocks * (iBlock + 1));
454template <
typename V,
typename T,
typename S>
458 uint32_t dstOffset = 0;
461 T* bufT =
reinterpret_cast<T*
>(
buf);
462 constexpr const int32_t
bufSize = GPUCA_WARP_SIZE;
463 constexpr const int32_t bufTSize =
bufSize *
sizeof(V) /
sizeof(T);
465 for (uint32_t
i = 0;
i < nEntries;
i++) {
467 uint32_t srcOffset = (srcOffsets[
i] * scaleBase1024 / 1024) + diff;
468 uint32_t srcSize = nums[
i] - diff;
470 if (dstAligned ==
nullptr) {
471 if (not isAlignedTo<V>(
dst)) {
472 size_t dsti =
reinterpret_cast<size_t>(
dst);
473 uint32_t
offset = (
alignof(V) - dsti %
alignof(V)) /
sizeof(T);
475 compressorMemcpyBasic(
dst,
src + srcOffset,
offset, nLanes, iLane);
479 if (isAlignedTo<V>(
dst)) {
480 dstAligned =
reinterpret_cast<V*
>(
dst);
483 while (srcPos < srcSize) {
484 uint32_t shmElemsLeft = bufTSize - shmPos;
485 uint32_t srcElemsLeft = srcSize - srcPos;
486 uint32_t
size = CAMath::Min(srcElemsLeft, shmElemsLeft);
487 compressorMemcpyBasic(bufT + shmPos,
src + srcOffset + srcPos,
size, nLanes, iLane);
492 if (shmPos >= bufTSize) {
493 compressorMemcpyBasic(dstAligned + dstOffset,
buf,
bufSize, nLanes, iLane);
501 compressorMemcpyBasic(
reinterpret_cast<T*
>(dstAligned + dstOffset), bufT, shmPos, nLanes, iLane);
508 uint32_t blockOffset = 0;
509 int32_t iThread = nLanes * iWarp + iLane;
510 int32_t nThreads = nLanes * nWarps;
511 uint32_t blockStart = work_group_broadcast(
start, 0);
512 for (uint32_t
i = iThread;
i < blockStart;
i += nThreads) {
513 blockOffset += nums[
i];
515 blockOffset = work_group_reduce_add(blockOffset);
518 for (uint32_t
i =
start + iLane;
i <
end;
i += nLanes) {
522 if (iWarp > -1 && iLane == nLanes - 1) {
523 smem.warpOffset[iWarp] =
offset;
526 offset = (iWarp <= 0) ? 0 : smem.warpOffset[iWarp - 1];
529 return offset + blockOffset;
538 int32_t nWarps = nThreads / GPUCA_WARP_SIZE;
539 int32_t iWarp = iThread / GPUCA_WARP_SIZE;
541 int32_t nLanes = GPUCA_WARP_SIZE;
542 int32_t iLane = iThread % GPUCA_WARP_SIZE;
547 uint32_t rowsPerWarp = (nRows + nWarps - 1) / nWarps;
548 uint32_t rowStart = rowsPerWarp * iWarp;
549 uint32_t rowEnd = CAMath::Min(nRows, rowStart + rowsPerWarp);
550 if (rowStart >= nRows) {
555 uint32_t rowsOffset = calculateWarpOffsets(smem, compressor.
mPtrs.
nSliceRowClusters, rowStart, rowEnd, nWarps, iWarp, nLanes, iLane);
571 for (uint32_t
i = sectorStart;
i <= sectorEnd &&
i < compressor.
NSECTORS;
i++) {
572 for (uint32_t
j = ((
i == sectorStart) ? sectorRowStart : 0);
j < ((
i == sectorEnd) ? sectorRowEnd :
GPUCA_ROW_COUNT);
j++) {
589 uint32_t trackStart = tracksPerWarp * iWarp;
596 uint32_t tracksOffset = calculateWarpOffsets(smem, compressor.
mPtrs.
nTrackClusters, trackStart, trackEnd, nWarps, iWarp, nLanes, iLane);
598 for (uint32_t
i = trackStart;
i < trackEnd;
i += nLanes) {
599 uint32_t nTrackClusters = 0;
600 uint32_t srcOffset = 0;
602 if (
i + iLane < trackEnd) {
606 smem.unbuffered.sizes[iWarp][iLane] = nTrackClusters;
607 smem.unbuffered.srcOffsets[iWarp][iLane] = srcOffset;
609 uint32_t elems = (
i + nLanes < trackEnd) ? nLanes : (trackEnd -
i);
611 for (uint32_t
j = 0;
j < elems;
j++) {
612 nTrackClusters = smem.unbuffered.sizes[iWarp][
j];
613 srcOffset = smem.unbuffered.srcOffsets[iWarp][
j];
614 uint32_t idx =
i +
j;
615 compressorMemcpy(compressor.
mOutput->
qTotA + tracksOffset, compressor.
mPtrs.
qTotA + srcOffset, nTrackClusters, nLanes, iLane);
616 compressorMemcpy(compressor.
mOutput->
qMaxA + tracksOffset, compressor.
mPtrs.
qMaxA + srcOffset, nTrackClusters, nLanes, iLane);
617 compressorMemcpy(compressor.
mOutput->
flagsA + tracksOffset, compressor.
mPtrs.
flagsA + srcOffset, nTrackClusters, nLanes, iLane);
622 compressorMemcpy(compressor.
mOutput->
rowDiffA + tracksOffset - idx, compressor.
mPtrs.
rowDiffA + srcOffset + 1, (nTrackClusters - 1), nLanes, iLane);
624 compressorMemcpy(compressor.
mOutput->
padResA + tracksOffset - idx, compressor.
mPtrs.
padResA + srcOffset + 1, (nTrackClusters - 1), nLanes, iLane);
625 compressorMemcpy(compressor.
mOutput->
timeResA + tracksOffset - idx, compressor.
mPtrs.
timeResA + srcOffset + 1, (nTrackClusters - 1), nLanes, iLane);
627 tracksOffset += nTrackClusters;
640 int32_t nWarps = nThreads / GPUCA_WARP_SIZE;
641 int32_t iWarp = iThread / GPUCA_WARP_SIZE;
643 int32_t nGlobalWarps = nWarps * nBlocks;
644 int32_t iGlobalWarp = nWarps * iBlock + iWarp;
646 int32_t nLanes = GPUCA_WARP_SIZE;
647 int32_t iLane = iThread % GPUCA_WARP_SIZE;
649 auto& input = compressor.
mPtrs;
653 uint32_t rowsPerWarp = (nRows + nGlobalWarps - 1) / nGlobalWarps;
654 uint32_t rowStart = rowsPerWarp * iGlobalWarp;
655 uint32_t rowEnd = CAMath::Min(nRows, rowStart + rowsPerWarp);
656 if (rowStart >= nRows) {
660 rowsPerWarp = rowEnd - rowStart;
662 uint32_t rowsOffset = calculateWarpOffsets(smem, input.nSliceRowClusters, rowStart, rowEnd, nWarps, iWarp, nLanes, iLane);
665 uint32_t tracksPerWarp = (nStoredTracks + nGlobalWarps - 1) / nGlobalWarps;
666 uint32_t trackStart = tracksPerWarp * iGlobalWarp;
667 uint32_t trackEnd = CAMath::Min(nStoredTracks, trackStart + tracksPerWarp);
668 if (trackStart >= nStoredTracks) {
672 tracksPerWarp = trackEnd - trackStart;
674 uint32_t tracksOffset = calculateWarpOffsets(smem, input.nTrackClusters, trackStart, trackEnd, nWarps, iWarp, nLanes, iLane);
686 const uint32_t* clusterOffsets = &
clusters->clusterOffset[0][0] + rowStart;
687 const uint32_t* nSectorRowClusters = input.nSliceRowClusters + rowStart;
689 auto*
buf = smem.getBuffer<V>(iWarp);
691 compressorMemcpyBuffered(
buf,
output->qTotU + rowsOffset, input.qTotU, nSectorRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0, compressor.
mMaxClusterFactorBase1024);
692 compressorMemcpyBuffered(
buf,
output->qMaxU + rowsOffset, input.qMaxU, nSectorRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0, compressor.
mMaxClusterFactorBase1024);
693 compressorMemcpyBuffered(
buf,
output->flagsU + rowsOffset, input.flagsU, nSectorRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0, compressor.
mMaxClusterFactorBase1024);
694 compressorMemcpyBuffered(
buf,
output->padDiffU + rowsOffset, input.padDiffU, nSectorRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0, compressor.
mMaxClusterFactorBase1024);
695 compressorMemcpyBuffered(
buf,
output->timeDiffU + rowsOffset, input.timeDiffU, nSectorRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0, compressor.
mMaxClusterFactorBase1024);
696 compressorMemcpyBuffered(
buf,
output->sigmaPadU + rowsOffset, input.sigmaPadU, nSectorRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0, compressor.
mMaxClusterFactorBase1024);
697 compressorMemcpyBuffered(
buf,
output->sigmaTimeU + rowsOffset, input.sigmaTimeU, nSectorRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0, compressor.
mMaxClusterFactorBase1024);
699 const uint16_t* nTrackClustersPtr = input.nTrackClusters + trackStart;
702 compressorMemcpyBuffered(
buf,
output->qTotA + tracksOffset, input.qTotA, nTrackClustersPtr, aClsFstIdx, tracksPerWarp, nLanes, iLane, 0);
703 compressorMemcpyBuffered(
buf,
output->qMaxA + tracksOffset, input.qMaxA, nTrackClustersPtr, aClsFstIdx, tracksPerWarp, nLanes, iLane, 0);
704 compressorMemcpyBuffered(
buf,
output->flagsA + tracksOffset, input.flagsA, nTrackClustersPtr, aClsFstIdx, tracksPerWarp, nLanes, iLane, 0);
705 compressorMemcpyBuffered(
buf,
output->sigmaPadA + tracksOffset, input.sigmaPadA, nTrackClustersPtr, aClsFstIdx, tracksPerWarp, nLanes, iLane, 0);
706 compressorMemcpyBuffered(
buf,
output->sigmaTimeA + tracksOffset, input.sigmaTimeA, nTrackClustersPtr, aClsFstIdx, tracksPerWarp, nLanes, iLane, 0);
709 uint32_t tracksOffsetDiff = tracksOffset - trackStart;
710 compressorMemcpyBuffered(
buf,
output->rowDiffA + tracksOffsetDiff, input.rowDiffA, nTrackClustersPtr, aClsFstIdx, tracksPerWarp, nLanes, iLane, 1);
711 compressorMemcpyBuffered(
buf,
output->sliceLegDiffA + tracksOffsetDiff, input.sliceLegDiffA, nTrackClustersPtr, aClsFstIdx, tracksPerWarp, nLanes, iLane, 1);
712 compressorMemcpyBuffered(
buf,
output->padResA + tracksOffsetDiff, input.padResA, nTrackClustersPtr, aClsFstIdx, tracksPerWarp, nLanes, iLane, 1);
713 compressorMemcpyBuffered(
buf,
output->timeResA + tracksOffsetDiff, input.timeResA, nTrackClustersPtr, aClsFstIdx, tracksPerWarp, nLanes, iLane, 1);
720 const auto& input = compressor.mPtrs;
721 auto*
output = compressor.mOutput;
723 const int32_t nWarps = nThreads / GPUCA_WARP_SIZE;
724 const int32_t iWarp = iThread / GPUCA_WARP_SIZE;
725 const int32_t nLanes = GPUCA_WARP_SIZE;
726 const int32_t iLane = iThread % GPUCA_WARP_SIZE;
730 compressorMemcpyBasic(
output->nSliceRowClusters, input.nSliceRowClusters, compressor.NSECTORS *
GPUCA_ROW_COUNT, nThreads, iThread);
731 compressorMemcpyBasic(
output->nTrackClusters, input.nTrackClusters, compressor.mMemory->nStoredTracks, nThreads, iThread);
732 compressorMemcpyBasic(
output->qPtA, input.qPtA, compressor.mMemory->nStoredTracks, nThreads, iThread);
733 compressorMemcpyBasic(
output->rowA, input.rowA, compressor.mMemory->nStoredTracks, nThreads, iThread);
734 compressorMemcpyBasic(
output->sliceA, input.sliceA, compressor.mMemory->nStoredTracks, nThreads, iThread);
735 compressorMemcpyBasic(
output->timeA, input.timeA, compressor.mMemory->nStoredTracks, nThreads, iThread);
736 compressorMemcpyBasic(
output->padA, input.padA, compressor.mMemory->nStoredTracks, nThreads, iThread);
737 }
else if (iBlock & 1) {
738 const uint32_t nGlobalWarps = nWarps * (nBlocks - 1) / 2;
739 const uint32_t iGlobalWarp = nWarps * (iBlock - 1) / 2 + iWarp;
742 uint32_t rowsPerWarp = (
nRows + nGlobalWarps - 1) / nGlobalWarps;
743 uint32_t rowStart = rowsPerWarp * iGlobalWarp;
744 uint32_t rowEnd = CAMath::Min(nRows, rowStart + rowsPerWarp);
745 if (rowStart >= nRows) {
749 rowsPerWarp = rowEnd - rowStart;
751 const uint32_t rowsOffset = calculateWarpOffsets(smem, input.nSliceRowClusters, rowStart, rowEnd, nWarps, iWarp, nLanes, iLane);
752 const uint32_t* clusterOffsets = &
clusters->clusterOffset[0][0] + rowStart;
753 const uint32_t* nSectorRowClusters = input.nSliceRowClusters + rowStart;
755 compressorMemcpyBuffered(
buf,
output->qTotU + rowsOffset, input.qTotU, nSectorRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0, compressor.mMaxClusterFactorBase1024);
756 compressorMemcpyBuffered(
buf,
output->qMaxU + rowsOffset, input.qMaxU, nSectorRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0, compressor.mMaxClusterFactorBase1024);
757 compressorMemcpyBuffered(
buf,
output->flagsU + rowsOffset, input.flagsU, nSectorRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0, compressor.mMaxClusterFactorBase1024);
758 compressorMemcpyBuffered(
buf,
output->padDiffU + rowsOffset, input.padDiffU, nSectorRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0, compressor.mMaxClusterFactorBase1024);
759 compressorMemcpyBuffered(
buf,
output->timeDiffU + rowsOffset, input.timeDiffU, nSectorRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0, compressor.mMaxClusterFactorBase1024);
760 compressorMemcpyBuffered(
buf,
output->sigmaPadU + rowsOffset, input.sigmaPadU, nSectorRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0, compressor.mMaxClusterFactorBase1024);
761 compressorMemcpyBuffered(
buf,
output->sigmaTimeU + rowsOffset, input.sigmaTimeU, nSectorRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0, compressor.mMaxClusterFactorBase1024);
763 const uint32_t nGlobalWarps = nWarps * (nBlocks - 1) / 2;
764 const uint32_t iGlobalWarp = nWarps * (iBlock / 2 - 1) + iWarp;
766 const uint32_t nStoredTracks = compressor.mMemory->nStoredTracks;
767 uint32_t tracksPerWarp = (nStoredTracks + nGlobalWarps - 1) / nGlobalWarps;
768 uint32_t trackStart = tracksPerWarp * iGlobalWarp;
769 uint32_t trackEnd = CAMath::Min(nStoredTracks, trackStart + tracksPerWarp);
770 if (trackStart >= nStoredTracks) {
774 tracksPerWarp = trackEnd - trackStart;
776 const uint32_t tracksOffset = calculateWarpOffsets(smem, input.nTrackClusters, trackStart, trackEnd, nWarps, iWarp, nLanes, iLane);
777 const uint16_t* nTrackClustersPtr = input.nTrackClusters + trackStart;
778 const uint32_t* aClsFstIdx = compressor.mAttachedClusterFirstIndex + trackStart;
780 compressorMemcpyBuffered(
buf,
output->qTotA + tracksOffset, input.qTotA, nTrackClustersPtr, aClsFstIdx, tracksPerWarp, nLanes, iLane, 0);
781 compressorMemcpyBuffered(
buf,
output->qMaxA + tracksOffset, input.qMaxA, nTrackClustersPtr, aClsFstIdx, tracksPerWarp, nLanes, iLane, 0);
782 compressorMemcpyBuffered(
buf,
output->flagsA + tracksOffset, input.flagsA, nTrackClustersPtr, aClsFstIdx, tracksPerWarp, nLanes, iLane, 0);
783 compressorMemcpyBuffered(
buf,
output->sigmaPadA + tracksOffset, input.sigmaPadA, nTrackClustersPtr, aClsFstIdx, tracksPerWarp, nLanes, iLane, 0);
784 compressorMemcpyBuffered(
buf,
output->sigmaTimeA + tracksOffset, input.sigmaTimeA, nTrackClustersPtr, aClsFstIdx, tracksPerWarp, nLanes, iLane, 0);
787 uint32_t tracksOffsetDiff = tracksOffset - trackStart;
788 compressorMemcpyBuffered(
buf,
output->rowDiffA + tracksOffsetDiff, input.rowDiffA, nTrackClustersPtr, aClsFstIdx, tracksPerWarp, nLanes, iLane, 1);
789 compressorMemcpyBuffered(
buf,
output->sliceLegDiffA + tracksOffsetDiff, input.sliceLegDiffA, nTrackClustersPtr, aClsFstIdx, tracksPerWarp, nLanes, iLane, 1);
790 compressorMemcpyBuffered(
buf,
output->padResA + tracksOffsetDiff, input.padResA, nTrackClustersPtr, aClsFstIdx, tracksPerWarp, nLanes, iLane, 1);
791 compressorMemcpyBuffered(
buf,
output->timeResA + tracksOffsetDiff, input.timeResA, nTrackClustersPtr, aClsFstIdx, tracksPerWarp, nLanes, iLane, 1);
798 gatherBuffered<Vec32>(nBlocks, nThreads, iBlock, iThread, smem, processors);
804 gatherBuffered<Vec64>(nBlocks, nThreads, iBlock, iThread, smem, processors);
810 gatherBuffered<Vec128>(nBlocks, nThreads, iBlock, iThread, smem, processors);
816 gatherMulti(nBlocks, nThreads, iBlock, iThread, smem, processors);
#define get_local_size(dim)
#define get_local_id(dim)
#define get_global_size(dim)
#define get_global_id(dim)
#define GPUCA_TPC_COMP_CHUNK_SIZE
#define GPUCA_GET_THREAD_COUNT(...)
GPUdii() void GPUTPCCompressionKernels
uint32_t * mAttachedClusterFirstIndex
o2::tpc::CompressedClusters * mOutput
static constexpr uint32_t NSECTORS
o2::tpc::CompressedClustersPtrs mPtrs
size_t mMaxClusterFactorBase1024
GLboolean GLboolean GLboolean b
typedef void(APIENTRYP PFNGLCULLFACEPROC)(GLenum mode)
GLboolean GLboolean GLboolean GLboolean a
GLenum GLuint GLenum GLsizei const GLchar * buf
GLdouble GLdouble GLdouble z
Global TPC definitions and constants.
GPUdi() T BetheBlochAleph(T bg
constexpr std::array< int, nLayers > nRows
a couple of static helper functions to create timestamp values for CCDB queries or override obsolete ...
static constexpr bool GetIsRejected(int32_t attach)
uint32_t nStoredAttachedClusters
const o2::tpc::ClusterNativeAccess * clustersNative
const uint32_t * mergedTrackHitAttachment
const GPUTPCGMMergedTrackHit * mergedTrackHits
const GPUTPCGMMergedTrack * mergedTracks
std::vector< std::byte > getBuffer(const char *filename)
std::vector< Cluster > clusters
for(int irof=0;irof< 1000;irof++)