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;
188 if (mClusterStatus[idx]) {
191 int32_t attach = ioPtrs.mergedTrackHitAttachment[
idx];
192 bool unattached = attach == 0;
203 auto& trk = ioPtrs.mergedTracks[
id];
204 if (CAMath::Abs(trk.GetParam().GetQPt() *
param.qptB5Scaler) >
param.rec.tpc.rejectQPtB5 || trk.MergedLooper()) {
218 uint32_t* sortBuffer = smem.sortBuffer;
222 const uint32_t idOffset =
clusters->clusterOffset[iSector][iRow];
225 if (iThread == nThreads - 1) {
228 uint32_t totalCount = 0;
233 const uint32_t nn = CAMath::nextMultipleOf<GPUCA_GET_THREAD_COUNT(GPUCA_LB_GPUTPCCompressionKernels_step1unattached)>(
clusters->nClusters[iSector][iRow]);
234 for (uint32_t
i = iThread;
i < nn + nThreads;
i += nThreads) {
235 const int32_t idx = idOffset +
i;
236 int32_t storeCluster =
i <
clusters->nClusters[iSector][iRow] && !compressor.rejectCluster(idx,
param, ioPtrs);
239 int32_t myIndex = work_group_scan_inclusive_add(storeCluster);
240 int32_t storeLater = -1;
243 sortBuffer[smem.nCount + myIndex - 1] =
i;
249 if (iThread == nThreads - 1) {
250 smem.nCount += myIndex;
259 if (idOffsetOut + totalCount +
count > idOffsetOutMax) {
260 if (iThread == nThreads - 1) {
261 compressor.raiseError(GPUErrors::ERROR_COMPRESSION_ROW_HIT_OVERFLOW, iSector * 1000 + iRow, idOffsetOut + totalCount +
count, idOffsetOutMax);
269#ifdef GPUCA_DETERMINISTIC_MODE
270 CAAlgo::sortInBlock(sortBuffer, sortBuffer +
count, GPUTPCCompressionKernels_Compare<GPUSettings::SortZPadTime>(
clusters->clusters[iSector][iRow]));
273 CAAlgo::sortInBlock(sortBuffer, sortBuffer +
count, GPUTPCCompressionKernels_Compare<GPUSettings::SortZPadTime>(
clusters->clusters[iSector][iRow]));
275 CAAlgo::sortInBlock(sortBuffer, sortBuffer +
count, GPUTPCCompressionKernels_Compare<GPUSettings::SortZTimePad>(
clusters->clusters[iSector][iRow]));
277 CAAlgo::sortInBlock(sortBuffer, sortBuffer +
count, GPUTPCCompressionKernels_Compare<GPUSettings::SortPad>(
clusters->clusters[iSector][iRow]));
279 CAAlgo::sortInBlock(sortBuffer, sortBuffer +
count, GPUTPCCompressionKernels_Compare<GPUSettings::SortTime>(
clusters->clusters[iSector][iRow]));
286 int32_t outidx = idOffsetOut + totalCount +
j;
289 int32_t preId =
j != 0 ? (int32_t)sortBuffer[
j - 1] : (totalCount != 0 ? (int32_t)smem.lastIndex : -1);
290 GPUTPCCompression_EncodeUnattached(
param.rec.tpc.compressionTypeMask, orgCl,
c.timeDiffU[outidx],
c.padDiffU[outidx], preId == -1 ?
nullptr : &
clusters->clusters[iSector][iRow][preId]);
292 uint16_t qtot = orgCl.
qTot, qmax = orgCl.
qMax;
295 compressor.truncateSignificantBitsChargeMax(qmax,
param);
296 compressor.truncateSignificantBitsCharge(qtot,
param);
297 compressor.truncateSignificantBitsWidth(sigmapad,
param);
298 compressor.truncateSignificantBitsWidth(sigmatime,
param);
300 c.qTotU[outidx] = qtot;
301 c.qMaxU[outidx] = qmax;
302 c.sigmaPadU[outidx] = sigmapad;
303 c.sigmaTimeU[outidx] = sigmatime;
304 c.flagsU[outidx] = orgCl.getFlags();
308 if (storeLater >= 0) {
309 sortBuffer[storeLater] =
i;
312 if (iThread == nThreads - 1 &&
count) {
313 smem.lastIndex = sortBuffer[
count - 1];
314 smem.nCount -=
count;
318 if (iThread == nThreads - 1) {
341 return buf128[iWarp];
344template <
typename T,
typename S>
347 if constexpr (
alignof(
S) >=
alignof(T)) {
348 static_cast<void>(
ptr);
351 return reinterpret_cast<size_t>(
ptr) %
alignof(T) == 0;
358 constexpr const int32_t vec128Elems = CpyVector<uint8_t, Vec128>::Size;
359 constexpr const int32_t vec64Elems = CpyVector<uint8_t, Vec64>::Size;
360 constexpr const int32_t vec32Elems = CpyVector<uint8_t, Vec32>::Size;
361 constexpr const int32_t vec16Elems = CpyVector<uint8_t, Vec16>::Size;
363 if (
size >= uint32_t(nThreads * vec128Elems)) {
364 compressorMemcpyVectorised<uint8_t, Vec128>(
dst,
src,
size, nThreads, iThread);
365 }
else if (
size >= uint32_t(nThreads * vec64Elems)) {
366 compressorMemcpyVectorised<uint8_t, Vec64>(
dst,
src,
size, nThreads, iThread);
367 }
else if (
size >= uint32_t(nThreads * vec32Elems)) {
368 compressorMemcpyVectorised<uint8_t, Vec32>(
dst,
src,
size, nThreads, iThread);
369 }
else if (
size >= uint32_t(nThreads * vec16Elems)) {
370 compressorMemcpyVectorised<uint8_t, Vec16>(
dst,
src,
size, nThreads, iThread);
372 compressorMemcpyBasic(
dst,
src,
size, nThreads, iThread);
379 constexpr const int32_t vec128Elems = CpyVector<uint16_t, Vec128>::Size;
380 constexpr const int32_t vec64Elems = CpyVector<uint16_t, Vec64>::Size;
381 constexpr const int32_t vec32Elems = CpyVector<uint16_t, Vec32>::Size;
383 if (
size >= uint32_t(nThreads * vec128Elems)) {
384 compressorMemcpyVectorised<uint16_t, Vec128>(
dst,
src,
size, nThreads, iThread);
385 }
else if (
size >= uint32_t(nThreads * vec64Elems)) {
386 compressorMemcpyVectorised<uint16_t, Vec64>(
dst,
src,
size, nThreads, iThread);
387 }
else if (
size >= uint32_t(nThreads * vec32Elems)) {
388 compressorMemcpyVectorised<uint16_t, Vec32>(
dst,
src,
size, nThreads, iThread);
390 compressorMemcpyBasic(
dst,
src,
size, nThreads, iThread);
397 constexpr const int32_t vec128Elems = CpyVector<uint32_t, Vec128>::Size;
398 constexpr const int32_t vec64Elems = CpyVector<uint32_t, Vec64>::Size;
400 if (
size >= uint32_t(nThreads * vec128Elems)) {
401 compressorMemcpyVectorised<uint32_t, Vec128>(
dst,
src,
size, nThreads, iThread);
402 }
else if (
size >= uint32_t(nThreads * vec64Elems)) {
403 compressorMemcpyVectorised<uint32_t, Vec64>(
dst,
src,
size, nThreads, iThread);
405 compressorMemcpyBasic(
dst,
src,
size, nThreads, iThread);
409template <
typename Scalar,
typename BaseVector>
412 if (not isAlignedTo<BaseVector>(
dst)) {
413 size_t dsti =
reinterpret_cast<size_t>(
dst);
414 int32_t
offset = (
alignof(BaseVector) - dsti %
alignof(BaseVector)) /
sizeof(Scalar);
415 compressorMemcpyBasic(
dst,
src,
offset, nThreads, iThread);
421 BaseVector*
GPUrestrict() dstAligned = reinterpret_cast<BaseVector*>(
dst);
423 using CpyVec = CpyVector<Scalar, BaseVector>;
424 uint32_t sizeAligned =
size / CpyVec::
Size;
426 if (isAlignedTo<BaseVector>(
src)) {
427 const BaseVector*
GPUrestrict() srcAligned = reinterpret_cast<const BaseVector*>(
src);
428 compressorMemcpyBasic(dstAligned, srcAligned, sizeAligned, nThreads, iThread);
430 for (uint32_t
i = iThread;
i < sizeAligned;
i += nThreads) {
432 for (uint32_t
j = 0;
j < CpyVec::Size;
j++) {
433 buf.elems[
j] =
src[
i * CpyVec::Size +
j];
435 dstAligned[
i] =
buf.all;
439 int32_t leftovers =
size % CpyVec::Size;
440 compressorMemcpyBasic(
dst +
size - leftovers,
src +
size - leftovers, leftovers, nThreads, iThread);
446 uint32_t
start = (
size + nBlocks - 1) / nBlocks * iBlock + iThread;
447 uint32_t
end = CAMath::Min(
size, (
size + nBlocks - 1) / nBlocks * (iBlock + 1));
453template <
typename V,
typename T,
typename S>
457 uint32_t dstOffset = 0;
460 T* bufT =
reinterpret_cast<T*
>(
buf);
461 constexpr const int32_t
bufSize = GPUCA_WARP_SIZE;
462 constexpr const int32_t bufTSize =
bufSize *
sizeof(V) /
sizeof(T);
464 for (uint32_t
i = 0;
i < nEntries;
i++) {
466 uint32_t srcOffset = (srcOffsets[
i] * scaleBase1024 / 1024) + diff;
467 uint32_t srcSize = nums[
i] - diff;
469 if (dstAligned ==
nullptr) {
470 if (not isAlignedTo<V>(
dst)) {
471 size_t dsti =
reinterpret_cast<size_t>(
dst);
472 uint32_t
offset = (
alignof(V) - dsti %
alignof(V)) /
sizeof(T);
474 compressorMemcpyBasic(
dst,
src + srcOffset,
offset, nLanes, iLane);
478 if (isAlignedTo<V>(
dst)) {
479 dstAligned =
reinterpret_cast<V*
>(
dst);
482 while (srcPos < srcSize) {
483 uint32_t shmElemsLeft = bufTSize - shmPos;
484 uint32_t srcElemsLeft = srcSize - srcPos;
485 uint32_t
size = CAMath::Min(srcElemsLeft, shmElemsLeft);
486 compressorMemcpyBasic(bufT + shmPos,
src + srcOffset + srcPos,
size, nLanes, iLane);
491 if (shmPos >= bufTSize) {
492 compressorMemcpyBasic(dstAligned + dstOffset,
buf,
bufSize, nLanes, iLane);
500 compressorMemcpyBasic(
reinterpret_cast<T*
>(dstAligned + dstOffset), bufT, shmPos, nLanes, iLane);
507 uint32_t blockOffset = 0;
508 int32_t iThread = nLanes * iWarp + iLane;
509 int32_t nThreads = nLanes * nWarps;
510 uint32_t blockStart = work_group_broadcast(
start, 0);
511 for (uint32_t
i = iThread;
i < blockStart;
i += nThreads) {
512 blockOffset += nums[
i];
514 blockOffset = work_group_reduce_add(blockOffset);
517 for (uint32_t
i =
start + iLane;
i <
end;
i += nLanes) {
521 if (iWarp > -1 && iLane == nLanes - 1) {
522 smem.warpOffset[iWarp] =
offset;
525 offset = (iWarp <= 0) ? 0 : smem.warpOffset[iWarp - 1];
528 return offset + blockOffset;
537 int32_t nWarps = nThreads / GPUCA_WARP_SIZE;
538 int32_t iWarp = iThread / GPUCA_WARP_SIZE;
540 int32_t nLanes = GPUCA_WARP_SIZE;
541 int32_t iLane = iThread % GPUCA_WARP_SIZE;
546 uint32_t rowsPerWarp = (nRows + nWarps - 1) / nWarps;
547 uint32_t rowStart = rowsPerWarp * iWarp;
548 uint32_t rowEnd = CAMath::Min(nRows, rowStart + rowsPerWarp);
549 if (rowStart >= nRows) {
554 uint32_t rowsOffset = calculateWarpOffsets(smem, compressor.
mPtrs.
nSliceRowClusters, rowStart, rowEnd, nWarps, iWarp, nLanes, iLane);
570 for (uint32_t
i = sectorStart;
i <= sectorEnd &&
i < compressor.
NSECTORS;
i++) {
571 for (uint32_t
j = ((
i == sectorStart) ? sectorRowStart : 0);
j < ((
i == sectorEnd) ? sectorRowEnd :
GPUCA_ROW_COUNT);
j++) {
588 uint32_t trackStart = tracksPerWarp * iWarp;
595 uint32_t tracksOffset = calculateWarpOffsets(smem, compressor.
mPtrs.
nTrackClusters, trackStart, trackEnd, nWarps, iWarp, nLanes, iLane);
597 for (uint32_t
i = trackStart;
i < trackEnd;
i += nLanes) {
598 uint32_t nTrackClusters = 0;
599 uint32_t srcOffset = 0;
601 if (
i + iLane < trackEnd) {
605 smem.unbuffered.sizes[iWarp][iLane] = nTrackClusters;
606 smem.unbuffered.srcOffsets[iWarp][iLane] = srcOffset;
608 uint32_t elems = (
i + nLanes < trackEnd) ? nLanes : (trackEnd -
i);
610 for (uint32_t
j = 0;
j < elems;
j++) {
611 nTrackClusters = smem.unbuffered.sizes[iWarp][
j];
612 srcOffset = smem.unbuffered.srcOffsets[iWarp][
j];
613 uint32_t idx =
i +
j;
614 compressorMemcpy(compressor.
mOutput->
qTotA + tracksOffset, compressor.
mPtrs.
qTotA + srcOffset, nTrackClusters, nLanes, iLane);
615 compressorMemcpy(compressor.
mOutput->
qMaxA + tracksOffset, compressor.
mPtrs.
qMaxA + srcOffset, nTrackClusters, nLanes, iLane);
616 compressorMemcpy(compressor.
mOutput->
flagsA + tracksOffset, compressor.
mPtrs.
flagsA + srcOffset, nTrackClusters, nLanes, iLane);
621 compressorMemcpy(compressor.
mOutput->
rowDiffA + tracksOffset - idx, compressor.
mPtrs.
rowDiffA + srcOffset + 1, (nTrackClusters - 1), nLanes, iLane);
623 compressorMemcpy(compressor.
mOutput->
padResA + tracksOffset - idx, compressor.
mPtrs.
padResA + srcOffset + 1, (nTrackClusters - 1), nLanes, iLane);
624 compressorMemcpy(compressor.
mOutput->
timeResA + tracksOffset - idx, compressor.
mPtrs.
timeResA + srcOffset + 1, (nTrackClusters - 1), nLanes, iLane);
626 tracksOffset += nTrackClusters;
639 int32_t nWarps = nThreads / GPUCA_WARP_SIZE;
640 int32_t iWarp = iThread / GPUCA_WARP_SIZE;
642 int32_t nGlobalWarps = nWarps * nBlocks;
643 int32_t iGlobalWarp = nWarps * iBlock + iWarp;
645 int32_t nLanes = GPUCA_WARP_SIZE;
646 int32_t iLane = iThread % GPUCA_WARP_SIZE;
648 auto& input = compressor.
mPtrs;
652 uint32_t rowsPerWarp = (nRows + nGlobalWarps - 1) / nGlobalWarps;
653 uint32_t rowStart = rowsPerWarp * iGlobalWarp;
654 uint32_t rowEnd = CAMath::Min(nRows, rowStart + rowsPerWarp);
655 if (rowStart >= nRows) {
659 rowsPerWarp = rowEnd - rowStart;
661 uint32_t rowsOffset = calculateWarpOffsets(smem, input.nSliceRowClusters, rowStart, rowEnd, nWarps, iWarp, nLanes, iLane);
664 uint32_t tracksPerWarp = (nStoredTracks + nGlobalWarps - 1) / nGlobalWarps;
665 uint32_t trackStart = tracksPerWarp * iGlobalWarp;
666 uint32_t trackEnd = CAMath::Min(nStoredTracks, trackStart + tracksPerWarp);
667 if (trackStart >= nStoredTracks) {
671 tracksPerWarp = trackEnd - trackStart;
673 uint32_t tracksOffset = calculateWarpOffsets(smem, input.nTrackClusters, trackStart, trackEnd, nWarps, iWarp, nLanes, iLane);
685 const uint32_t* clusterOffsets = &
clusters->clusterOffset[0][0] + rowStart;
686 const uint32_t* nSectorRowClusters = input.nSliceRowClusters + rowStart;
688 auto*
buf = smem.getBuffer<V>(iWarp);
690 compressorMemcpyBuffered(
buf,
output->qTotU + rowsOffset, input.qTotU, nSectorRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0, compressor.
mMaxClusterFactorBase1024);
691 compressorMemcpyBuffered(
buf,
output->qMaxU + rowsOffset, input.qMaxU, nSectorRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0, compressor.
mMaxClusterFactorBase1024);
692 compressorMemcpyBuffered(
buf,
output->flagsU + rowsOffset, input.flagsU, nSectorRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0, compressor.
mMaxClusterFactorBase1024);
693 compressorMemcpyBuffered(
buf,
output->padDiffU + rowsOffset, input.padDiffU, nSectorRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0, compressor.
mMaxClusterFactorBase1024);
694 compressorMemcpyBuffered(
buf,
output->timeDiffU + rowsOffset, input.timeDiffU, nSectorRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0, compressor.
mMaxClusterFactorBase1024);
695 compressorMemcpyBuffered(
buf,
output->sigmaPadU + rowsOffset, input.sigmaPadU, nSectorRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0, compressor.
mMaxClusterFactorBase1024);
696 compressorMemcpyBuffered(
buf,
output->sigmaTimeU + rowsOffset, input.sigmaTimeU, nSectorRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0, compressor.
mMaxClusterFactorBase1024);
698 const uint16_t* nTrackClustersPtr = input.nTrackClusters + trackStart;
701 compressorMemcpyBuffered(
buf,
output->qTotA + tracksOffset, input.qTotA, nTrackClustersPtr, aClsFstIdx, tracksPerWarp, nLanes, iLane, 0);
702 compressorMemcpyBuffered(
buf,
output->qMaxA + tracksOffset, input.qMaxA, nTrackClustersPtr, aClsFstIdx, tracksPerWarp, nLanes, iLane, 0);
703 compressorMemcpyBuffered(
buf,
output->flagsA + tracksOffset, input.flagsA, nTrackClustersPtr, aClsFstIdx, tracksPerWarp, nLanes, iLane, 0);
704 compressorMemcpyBuffered(
buf,
output->sigmaPadA + tracksOffset, input.sigmaPadA, nTrackClustersPtr, aClsFstIdx, tracksPerWarp, nLanes, iLane, 0);
705 compressorMemcpyBuffered(
buf,
output->sigmaTimeA + tracksOffset, input.sigmaTimeA, nTrackClustersPtr, aClsFstIdx, tracksPerWarp, nLanes, iLane, 0);
708 uint32_t tracksOffsetDiff = tracksOffset - trackStart;
709 compressorMemcpyBuffered(
buf,
output->rowDiffA + tracksOffsetDiff, input.rowDiffA, nTrackClustersPtr, aClsFstIdx, tracksPerWarp, nLanes, iLane, 1);
710 compressorMemcpyBuffered(
buf,
output->sliceLegDiffA + tracksOffsetDiff, input.sliceLegDiffA, nTrackClustersPtr, aClsFstIdx, tracksPerWarp, nLanes, iLane, 1);
711 compressorMemcpyBuffered(
buf,
output->padResA + tracksOffsetDiff, input.padResA, nTrackClustersPtr, aClsFstIdx, tracksPerWarp, nLanes, iLane, 1);
712 compressorMemcpyBuffered(
buf,
output->timeResA + tracksOffsetDiff, input.timeResA, nTrackClustersPtr, aClsFstIdx, tracksPerWarp, nLanes, iLane, 1);
719 const auto& input = compressor.mPtrs;
720 auto*
output = compressor.mOutput;
722 const int32_t nWarps = nThreads / GPUCA_WARP_SIZE;
723 const int32_t iWarp = iThread / GPUCA_WARP_SIZE;
724 const int32_t nLanes = GPUCA_WARP_SIZE;
725 const int32_t iLane = iThread % GPUCA_WARP_SIZE;
729 compressorMemcpyBasic(
output->nSliceRowClusters, input.nSliceRowClusters, compressor.NSECTORS *
GPUCA_ROW_COUNT, nThreads, iThread);
730 compressorMemcpyBasic(
output->nTrackClusters, input.nTrackClusters, compressor.mMemory->nStoredTracks, nThreads, iThread);
731 compressorMemcpyBasic(
output->qPtA, input.qPtA, compressor.mMemory->nStoredTracks, nThreads, iThread);
732 compressorMemcpyBasic(
output->rowA, input.rowA, compressor.mMemory->nStoredTracks, nThreads, iThread);
733 compressorMemcpyBasic(
output->sliceA, input.sliceA, compressor.mMemory->nStoredTracks, nThreads, iThread);
734 compressorMemcpyBasic(
output->timeA, input.timeA, compressor.mMemory->nStoredTracks, nThreads, iThread);
735 compressorMemcpyBasic(
output->padA, input.padA, compressor.mMemory->nStoredTracks, nThreads, iThread);
736 }
else if (iBlock & 1) {
737 const uint32_t nGlobalWarps = nWarps * (nBlocks - 1) / 2;
738 const uint32_t iGlobalWarp = nWarps * (iBlock - 1) / 2 + iWarp;
741 uint32_t rowsPerWarp = (
nRows + nGlobalWarps - 1) / nGlobalWarps;
742 uint32_t rowStart = rowsPerWarp * iGlobalWarp;
743 uint32_t rowEnd = CAMath::Min(nRows, rowStart + rowsPerWarp);
744 if (rowStart >= nRows) {
748 rowsPerWarp = rowEnd - rowStart;
750 const uint32_t rowsOffset = calculateWarpOffsets(smem, input.nSliceRowClusters, rowStart, rowEnd, nWarps, iWarp, nLanes, iLane);
751 const uint32_t* clusterOffsets = &
clusters->clusterOffset[0][0] + rowStart;
752 const uint32_t* nSectorRowClusters = input.nSliceRowClusters + rowStart;
754 compressorMemcpyBuffered(
buf,
output->qTotU + rowsOffset, input.qTotU, nSectorRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0, compressor.mMaxClusterFactorBase1024);
755 compressorMemcpyBuffered(
buf,
output->qMaxU + rowsOffset, input.qMaxU, nSectorRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0, compressor.mMaxClusterFactorBase1024);
756 compressorMemcpyBuffered(
buf,
output->flagsU + rowsOffset, input.flagsU, nSectorRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0, compressor.mMaxClusterFactorBase1024);
757 compressorMemcpyBuffered(
buf,
output->padDiffU + rowsOffset, input.padDiffU, nSectorRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0, compressor.mMaxClusterFactorBase1024);
758 compressorMemcpyBuffered(
buf,
output->timeDiffU + rowsOffset, input.timeDiffU, nSectorRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0, compressor.mMaxClusterFactorBase1024);
759 compressorMemcpyBuffered(
buf,
output->sigmaPadU + rowsOffset, input.sigmaPadU, nSectorRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0, compressor.mMaxClusterFactorBase1024);
760 compressorMemcpyBuffered(
buf,
output->sigmaTimeU + rowsOffset, input.sigmaTimeU, nSectorRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0, compressor.mMaxClusterFactorBase1024);
762 const uint32_t nGlobalWarps = nWarps * (nBlocks - 1) / 2;
763 const uint32_t iGlobalWarp = nWarps * (iBlock / 2 - 1) + iWarp;
765 const uint32_t nStoredTracks = compressor.mMemory->nStoredTracks;
766 uint32_t tracksPerWarp = (nStoredTracks + nGlobalWarps - 1) / nGlobalWarps;
767 uint32_t trackStart = tracksPerWarp * iGlobalWarp;
768 uint32_t trackEnd = CAMath::Min(nStoredTracks, trackStart + tracksPerWarp);
769 if (trackStart >= nStoredTracks) {
773 tracksPerWarp = trackEnd - trackStart;
775 const uint32_t tracksOffset = calculateWarpOffsets(smem, input.nTrackClusters, trackStart, trackEnd, nWarps, iWarp, nLanes, iLane);
776 const uint16_t* nTrackClustersPtr = input.nTrackClusters + trackStart;
777 const uint32_t* aClsFstIdx = compressor.mAttachedClusterFirstIndex + trackStart;
779 compressorMemcpyBuffered(
buf,
output->qTotA + tracksOffset, input.qTotA, nTrackClustersPtr, aClsFstIdx, tracksPerWarp, nLanes, iLane, 0);
780 compressorMemcpyBuffered(
buf,
output->qMaxA + tracksOffset, input.qMaxA, nTrackClustersPtr, aClsFstIdx, tracksPerWarp, nLanes, iLane, 0);
781 compressorMemcpyBuffered(
buf,
output->flagsA + tracksOffset, input.flagsA, nTrackClustersPtr, aClsFstIdx, tracksPerWarp, nLanes, iLane, 0);
782 compressorMemcpyBuffered(
buf,
output->sigmaPadA + tracksOffset, input.sigmaPadA, nTrackClustersPtr, aClsFstIdx, tracksPerWarp, nLanes, iLane, 0);
783 compressorMemcpyBuffered(
buf,
output->sigmaTimeA + tracksOffset, input.sigmaTimeA, nTrackClustersPtr, aClsFstIdx, tracksPerWarp, nLanes, iLane, 0);
786 uint32_t tracksOffsetDiff = tracksOffset - trackStart;
787 compressorMemcpyBuffered(
buf,
output->rowDiffA + tracksOffsetDiff, input.rowDiffA, nTrackClustersPtr, aClsFstIdx, tracksPerWarp, nLanes, iLane, 1);
788 compressorMemcpyBuffered(
buf,
output->sliceLegDiffA + tracksOffsetDiff, input.sliceLegDiffA, nTrackClustersPtr, aClsFstIdx, tracksPerWarp, nLanes, iLane, 1);
789 compressorMemcpyBuffered(
buf,
output->padResA + tracksOffsetDiff, input.padResA, nTrackClustersPtr, aClsFstIdx, tracksPerWarp, nLanes, iLane, 1);
790 compressorMemcpyBuffered(
buf,
output->timeResA + tracksOffsetDiff, input.timeResA, nTrackClustersPtr, aClsFstIdx, tracksPerWarp, nLanes, iLane, 1);
797 gatherBuffered<Vec32>(nBlocks, nThreads, iBlock, iThread, smem, processors);
803 gatherBuffered<Vec64>(nBlocks, nThreads, iBlock, iThread, smem, processors);
809 gatherBuffered<Vec128>(nBlocks, nThreads, iBlock, iThread, smem, processors);
815 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
uint32_t nStoredUnattachedClusters
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