22#include "GPUTPCCompressionKernels.inc"
43 bool rejectTrk = CAMath::Abs(trk.GetParam().GetQPt() * processors.param.qptB5Scaler) > processors.param.rec.tpc.rejectQPtB5 || trk.MergedLooper();
44 uint32_t nClustersStored = 0;
46 uint8_t lastRow = 0, lastSector = 0;
49 for (int32_t k = trk.NClusters() - 1; k >= 0; k--) {
55 int32_t hitId = hit.
num;
71 float x = geo.Row2X(hit.
row);
72 float y = track.LinearPad2Y(hit.
sector, orgCl.getPad(), geo.PadWidth(hit.
row), geo.NPads(hit.
row));
73 float z = geo.LinearTime2Z(hit.
sector, orgCl.getTime());
74 if (nClustersStored) {
78 if (lastLeg != hit.
leg && track.Mirror()) {
81 if (track.Propagate(geo.Row2X(hit.
row),
param.SectorParam[hit.
sector].Alpha)) {
88 int32_t cidx = trk.FirstClusterRef() + nClustersStored++;
89 if (nClustersStored == 1) {
90 uint8_t qpt = fabs(trk.GetParam().GetQPt()) < 20.f ? (trk.GetParam().GetQPt() * (127.f / 20.f) + 127.5f) : (trk.GetParam().GetQPt() > 0 ? 254 : 0);
97 c.qPtA[myTrack] = qpt;
98 c.rowA[myTrack] = hit.
row;
100 c.timeA[myTrack] = orgCl.getTimePacked();
104 uint32_t sector = hit.
sector;
111 if (lastSector > sector) {
114 sector -= lastSector;
116 c.rowDiffA[cidx] =
row;
117 c.sliceLegDiffA[cidx] = (hit.
leg == lastLeg ? 0 : compressor.
NSECTORS) + sector;
118 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))));
119 c.padResA[cidx] = orgCl.
padPacked - orgCl.packPad(pad);
120 float time = CAMath::Max(0.f, geo.LinearZ2Time(hit.
sector, track.Z() + zOffset));
121 c.timeResA[cidx] = (orgCl.getTimePacked() - orgCl.packTime(
time)) & 0xFFFFFF;
124 uint16_t qtot = orgCl.
qTot, qmax = orgCl.
qMax;
127 compressor.truncateSignificantBitsChargeMax(qmax,
param);
128 compressor.truncateSignificantBitsCharge(qtot,
param);
129 compressor.truncateSignificantBitsWidth(sigmapad,
param);
130 compressor.truncateSignificantBitsWidth(sigmatime,
param);
132 c.qTotA[cidx] = qtot;
133 c.qMaxA[cidx] = qmax;
134 c.sigmaPadA[cidx] = sigmapad;
135 c.sigmaTimeA[cidx] = sigmatime;
136 c.flagsA[cidx] = orgCl.getFlags();
137 if (k && track.Filter(
y,
z - zOffset, hit.
row)) {
143 if (nClustersStored) {
145 c.nTrackClusters[myTrack] = nClustersStored;
153 return mClsPtr[
a].getTimePacked() < mClsPtr[
b].getTimePacked();
159 return mClsPtr[
a].padPacked < mClsPtr[
b].padPacked;
165 if (mClsPtr[
a].getTimePacked() >> 3 == mClsPtr[
b].getTimePacked() >> 3) {
166 return mClsPtr[
a].padPacked < mClsPtr[
b].padPacked;
168 return mClsPtr[
a].getTimePacked() < mClsPtr[
b].getTimePacked();
174 if (mClsPtr[
a].padPacked >> 3 == mClsPtr[
b].padPacked >> 3) {
175 return mClsPtr[
a].getTimePacked() < mClsPtr[
b].getTimePacked();
177 return mClsPtr[
a].padPacked < mClsPtr[
b].padPacked;
187 uint32_t* sortBuffer = smem.sortBuffer;
191 const uint32_t idOffset =
clusters->clusterOffset[iSector][iRow];
192 const uint32_t idOffsetOut =
clusters->clusterOffset[iSector][iRow] * compressor.mMaxClusterFactorBase1024 / 1024;
193 const uint32_t idOffsetOutMax = ((
const uint32_t*)
clusters->clusterOffset[iSector])[iRow + 1] * compressor.mMaxClusterFactorBase1024 / 1024;
194 if (iThread == nThreads - 1) {
197 uint32_t totalCount = 0;
203 for (uint32_t
i = iThread;
i < nn + nThreads;
i += nThreads) {
204 const int32_t
idx = idOffset +
i;
207 if (
i >=
clusters->nClusters[iSector][iRow]) {
210 if (compressor.mClusterStatus[idx]) {
213 int32_t attach = ioPtrs.mergedTrackHitAttachment[
idx];
214 bool unattached = attach == 0;
225 auto& trk = ioPtrs.mergedTracks[
id];
226 if (CAMath::Abs(trk.GetParam().GetQPt() * processors.param.qptB5Scaler) > processors.param.rec.tpc.rejectQPtB5 || trk.MergedLooper()) {
234 int32_t myIndex = work_group_scan_inclusive_add(cidx);
235 int32_t storeLater = -1;
238 sortBuffer[smem.nCount + myIndex - 1] =
i;
244 if (iThread == nThreads - 1) {
245 smem.nCount += myIndex;
254 if (idOffsetOut + totalCount +
count > idOffsetOutMax) {
255 if (iThread == nThreads - 1) {
256 compressor.raiseError(GPUErrors::ERROR_COMPRESSION_ROW_HIT_OVERFLOW, iSector * 1000 + iRow, idOffsetOut + totalCount +
count, idOffsetOutMax);
265 CAAlgo::sortInBlock(sortBuffer, sortBuffer +
count, GPUTPCCompressionKernels_Compare<GPUSettings::SortZPadTime>(
clusters->clusters[iSector][iRow]));
267 CAAlgo::sortInBlock(sortBuffer, sortBuffer +
count, GPUTPCCompressionKernels_Compare<GPUSettings::SortZTimePad>(
clusters->clusters[iSector][iRow]));
269 CAAlgo::sortInBlock(sortBuffer, sortBuffer +
count, GPUTPCCompressionKernels_Compare<GPUSettings::SortPad>(
clusters->clusters[iSector][iRow]));
271 CAAlgo::sortInBlock(sortBuffer, sortBuffer +
count, GPUTPCCompressionKernels_Compare<GPUSettings::SortTime>(
clusters->clusters[iSector][iRow]));
277 int32_t outidx = idOffsetOut + totalCount +
j;
280 int32_t preId =
j != 0 ? (int32_t)sortBuffer[
j - 1] : (totalCount != 0 ? (int32_t)smem.lastIndex : -1);
281 GPUTPCCompression_EncodeUnattached(
param.
rec.tpc.compressionTypeMask, orgCl,
c.timeDiffU[outidx],
c.padDiffU[outidx], preId == -1 ?
nullptr : &
clusters->
clusters[iSector][iRow][preId]);
283 uint16_t qtot = orgCl.qTot, qmax = orgCl.qMax;
284 uint8_t sigmapad = orgCl.sigmaPadPacked, sigmatime = orgCl.sigmaTimePacked;
286 compressor.truncateSignificantBitsChargeMax(qmax,
param);
287 compressor.truncateSignificantBitsCharge(qtot,
param);
288 compressor.truncateSignificantBitsWidth(sigmapad,
param);
289 compressor.truncateSignificantBitsWidth(sigmatime,
param);
291 c.qTotU[outidx] = qtot;
292 c.qMaxU[outidx] = qmax;
293 c.sigmaPadU[outidx] = sigmapad;
294 c.sigmaTimeU[outidx] = sigmatime;
295 c.flagsU[outidx] = orgCl.getFlags();
299 if (storeLater >= 0) {
300 sortBuffer[storeLater] =
i;
303 if (iThread == nThreads - 1 &&
count) {
304 smem.lastIndex = sortBuffer[
count - 1];
305 smem.nCount -=
count;
309 if (iThread == nThreads - 1) {
311 CAMath::AtomicAdd(&compressor.mMemory->nStoredUnattachedClusters, totalCount);
332 return buf128[iWarp];
335template <
typename T,
typename S>
338 if constexpr (
alignof(
S) >=
alignof(T)) {
339 static_cast<void>(
ptr);
342 return reinterpret_cast<size_t>(
ptr) %
alignof(T) == 0;
349 constexpr const int32_t vec128Elems = CpyVector<uint8_t, Vec128>::Size;
350 constexpr const int32_t vec64Elems = CpyVector<uint8_t, Vec64>::Size;
351 constexpr const int32_t vec32Elems = CpyVector<uint8_t, Vec32>::Size;
352 constexpr const int32_t vec16Elems = CpyVector<uint8_t, Vec16>::Size;
354 if (
size >= uint32_t(nThreads * vec128Elems)) {
355 compressorMemcpyVectorised<uint8_t, Vec128>(
dst,
src,
size, nThreads, iThread);
356 }
else if (
size >= uint32_t(nThreads * vec64Elems)) {
357 compressorMemcpyVectorised<uint8_t, Vec64>(
dst,
src,
size, nThreads, iThread);
358 }
else if (
size >= uint32_t(nThreads * vec32Elems)) {
359 compressorMemcpyVectorised<uint8_t, Vec32>(
dst,
src,
size, nThreads, iThread);
360 }
else if (
size >= uint32_t(nThreads * vec16Elems)) {
361 compressorMemcpyVectorised<uint8_t, Vec16>(
dst,
src,
size, nThreads, iThread);
363 compressorMemcpyBasic(
dst,
src,
size, nThreads, iThread);
370 constexpr const int32_t vec128Elems = CpyVector<uint16_t, Vec128>::Size;
371 constexpr const int32_t vec64Elems = CpyVector<uint16_t, Vec64>::Size;
372 constexpr const int32_t vec32Elems = CpyVector<uint16_t, Vec32>::Size;
374 if (
size >= uint32_t(nThreads * vec128Elems)) {
375 compressorMemcpyVectorised<uint16_t, Vec128>(
dst,
src,
size, nThreads, iThread);
376 }
else if (
size >= uint32_t(nThreads * vec64Elems)) {
377 compressorMemcpyVectorised<uint16_t, Vec64>(
dst,
src,
size, nThreads, iThread);
378 }
else if (
size >= uint32_t(nThreads * vec32Elems)) {
379 compressorMemcpyVectorised<uint16_t, Vec32>(
dst,
src,
size, nThreads, iThread);
381 compressorMemcpyBasic(
dst,
src,
size, nThreads, iThread);
388 constexpr const int32_t vec128Elems = CpyVector<uint32_t, Vec128>::Size;
389 constexpr const int32_t vec64Elems = CpyVector<uint32_t, Vec64>::Size;
391 if (
size >= uint32_t(nThreads * vec128Elems)) {
392 compressorMemcpyVectorised<uint32_t, Vec128>(
dst,
src,
size, nThreads, iThread);
393 }
else if (
size >= uint32_t(nThreads * vec64Elems)) {
394 compressorMemcpyVectorised<uint32_t, Vec64>(
dst,
src,
size, nThreads, iThread);
396 compressorMemcpyBasic(
dst,
src,
size, nThreads, iThread);
400template <
typename Scalar,
typename BaseVector>
403 if (not isAlignedTo<BaseVector>(
dst)) {
404 size_t dsti =
reinterpret_cast<size_t>(
dst);
405 int32_t
offset = (
alignof(BaseVector) - dsti %
alignof(BaseVector)) /
sizeof(Scalar);
406 compressorMemcpyBasic(
dst,
src,
offset, nThreads, iThread);
412 BaseVector*
GPUrestrict() dstAligned = reinterpret_cast<BaseVector*>(
dst);
414 using CpyVec = CpyVector<Scalar, BaseVector>;
415 uint32_t sizeAligned =
size / CpyVec::
Size;
417 if (isAlignedTo<BaseVector>(
src)) {
418 const BaseVector*
GPUrestrict() srcAligned = reinterpret_cast<const BaseVector*>(
src);
419 compressorMemcpyBasic(dstAligned, srcAligned, sizeAligned, nThreads, iThread);
421 for (uint32_t
i = iThread;
i < sizeAligned;
i += nThreads) {
423 for (uint32_t
j = 0;
j < CpyVec::Size;
j++) {
424 buf.elems[
j] =
src[
i * CpyVec::Size +
j];
426 dstAligned[
i] =
buf.all;
430 int32_t leftovers =
size % CpyVec::Size;
431 compressorMemcpyBasic(
dst +
size - leftovers,
src +
size - leftovers, leftovers, nThreads, iThread);
437 uint32_t
start = (
size + nBlocks - 1) / nBlocks * iBlock + iThread;
438 uint32_t
end = CAMath::Min(
size, (
size + nBlocks - 1) / nBlocks * (iBlock + 1));
444template <
typename V,
typename T,
typename S>
448 uint32_t dstOffset = 0;
451 T* bufT =
reinterpret_cast<T*
>(
buf);
452 constexpr const int32_t
bufSize = GPUCA_WARP_SIZE;
453 constexpr const int32_t bufTSize =
bufSize *
sizeof(V) /
sizeof(T);
455 for (uint32_t
i = 0;
i < nEntries;
i++) {
457 uint32_t srcOffset = (srcOffsets[
i] * scaleBase1024 / 1024) + diff;
458 uint32_t srcSize = nums[
i] - diff;
460 if (dstAligned ==
nullptr) {
461 if (not isAlignedTo<V>(
dst)) {
462 size_t dsti =
reinterpret_cast<size_t>(
dst);
463 uint32_t
offset = (
alignof(V) - dsti %
alignof(V)) /
sizeof(T);
465 compressorMemcpyBasic(
dst,
src + srcOffset,
offset, nLanes, iLane);
469 if (isAlignedTo<V>(
dst)) {
470 dstAligned =
reinterpret_cast<V*
>(
dst);
473 while (srcPos < srcSize) {
474 uint32_t shmElemsLeft = bufTSize - shmPos;
475 uint32_t srcElemsLeft = srcSize - srcPos;
476 uint32_t
size = CAMath::Min(srcElemsLeft, shmElemsLeft);
477 compressorMemcpyBasic(bufT + shmPos,
src + srcOffset + srcPos,
size, nLanes, iLane);
482 if (shmPos >= bufTSize) {
483 compressorMemcpyBasic(dstAligned + dstOffset,
buf,
bufSize, nLanes, iLane);
491 compressorMemcpyBasic(
reinterpret_cast<T*
>(dstAligned + dstOffset), bufT, shmPos, nLanes, iLane);
498 uint32_t blockOffset = 0;
499 int32_t iThread = nLanes * iWarp + iLane;
500 int32_t nThreads = nLanes * nWarps;
501 uint32_t blockStart = work_group_broadcast(
start, 0);
502 for (uint32_t
i = iThread;
i < blockStart;
i += nThreads) {
503 blockOffset += nums[
i];
505 blockOffset = work_group_reduce_add(blockOffset);
508 for (uint32_t
i =
start + iLane;
i <
end;
i += nLanes) {
512 if (iWarp > -1 && iLane == nLanes - 1) {
513 smem.warpOffset[iWarp] =
offset;
516 offset = (iWarp <= 0) ? 0 : smem.warpOffset[iWarp - 1];
519 return offset + blockOffset;
528 int32_t nWarps = nThreads / GPUCA_WARP_SIZE;
529 int32_t iWarp = iThread / GPUCA_WARP_SIZE;
531 int32_t nLanes = GPUCA_WARP_SIZE;
532 int32_t iLane = iThread % GPUCA_WARP_SIZE;
537 uint32_t rowsPerWarp = (nRows + nWarps - 1) / nWarps;
538 uint32_t rowStart = rowsPerWarp * iWarp;
539 uint32_t rowEnd = CAMath::Min(nRows, rowStart + rowsPerWarp);
540 if (rowStart >= nRows) {
545 uint32_t rowsOffset = calculateWarpOffsets(smem, compressor.
mPtrs.
nSliceRowClusters, rowStart, rowEnd, nWarps, iWarp, nLanes, iLane);
561 for (uint32_t
i = sectorStart;
i <= sectorEnd &&
i < compressor.
NSECTORS;
i++) {
562 for (uint32_t
j = ((
i == sectorStart) ? sectorRowStart : 0);
j < ((
i == sectorEnd) ? sectorRowEnd :
GPUCA_ROW_COUNT);
j++) {
579 uint32_t trackStart = tracksPerWarp * iWarp;
586 uint32_t tracksOffset = calculateWarpOffsets(smem, compressor.
mPtrs.
nTrackClusters, trackStart, trackEnd, nWarps, iWarp, nLanes, iLane);
588 for (uint32_t
i = trackStart;
i < trackEnd;
i += nLanes) {
589 uint32_t nTrackClusters = 0;
590 uint32_t srcOffset = 0;
592 if (
i + iLane < trackEnd) {
596 smem.unbuffered.sizes[iWarp][iLane] = nTrackClusters;
597 smem.unbuffered.srcOffsets[iWarp][iLane] = srcOffset;
599 uint32_t elems = (
i + nLanes < trackEnd) ? nLanes : (trackEnd -
i);
601 for (uint32_t
j = 0;
j < elems;
j++) {
602 nTrackClusters = smem.unbuffered.sizes[iWarp][
j];
603 srcOffset = smem.unbuffered.srcOffsets[iWarp][
j];
604 uint32_t idx =
i +
j;
605 compressorMemcpy(compressor.
mOutput->
qTotA + tracksOffset, compressor.
mPtrs.
qTotA + srcOffset, nTrackClusters, nLanes, iLane);
606 compressorMemcpy(compressor.
mOutput->
qMaxA + tracksOffset, compressor.
mPtrs.
qMaxA + srcOffset, nTrackClusters, nLanes, iLane);
607 compressorMemcpy(compressor.
mOutput->
flagsA + tracksOffset, compressor.
mPtrs.
flagsA + srcOffset, nTrackClusters, nLanes, iLane);
612 compressorMemcpy(compressor.
mOutput->
rowDiffA + tracksOffset - idx, compressor.
mPtrs.
rowDiffA + srcOffset + 1, (nTrackClusters - 1), nLanes, iLane);
614 compressorMemcpy(compressor.
mOutput->
padResA + tracksOffset - idx, compressor.
mPtrs.
padResA + srcOffset + 1, (nTrackClusters - 1), nLanes, iLane);
615 compressorMemcpy(compressor.
mOutput->
timeResA + tracksOffset - idx, compressor.
mPtrs.
timeResA + srcOffset + 1, (nTrackClusters - 1), nLanes, iLane);
617 tracksOffset += nTrackClusters;
630 int32_t nWarps = nThreads / GPUCA_WARP_SIZE;
631 int32_t iWarp = iThread / GPUCA_WARP_SIZE;
633 int32_t nGlobalWarps = nWarps * nBlocks;
634 int32_t iGlobalWarp = nWarps * iBlock + iWarp;
636 int32_t nLanes = GPUCA_WARP_SIZE;
637 int32_t iLane = iThread % GPUCA_WARP_SIZE;
639 auto& input = compressor.
mPtrs;
643 uint32_t rowsPerWarp = (nRows + nGlobalWarps - 1) / nGlobalWarps;
644 uint32_t rowStart = rowsPerWarp * iGlobalWarp;
645 uint32_t rowEnd = CAMath::Min(nRows, rowStart + rowsPerWarp);
646 if (rowStart >= nRows) {
650 rowsPerWarp = rowEnd - rowStart;
652 uint32_t rowsOffset = calculateWarpOffsets(smem, input.nSliceRowClusters, rowStart, rowEnd, nWarps, iWarp, nLanes, iLane);
655 uint32_t tracksPerWarp = (nStoredTracks + nGlobalWarps - 1) / nGlobalWarps;
656 uint32_t trackStart = tracksPerWarp * iGlobalWarp;
657 uint32_t trackEnd = CAMath::Min(nStoredTracks, trackStart + tracksPerWarp);
658 if (trackStart >= nStoredTracks) {
662 tracksPerWarp = trackEnd - trackStart;
664 uint32_t tracksOffset = calculateWarpOffsets(smem, input.nTrackClusters, trackStart, trackEnd, nWarps, iWarp, nLanes, iLane);
676 const uint32_t* clusterOffsets = &
clusters->clusterOffset[0][0] + rowStart;
677 const uint32_t* nSectorRowClusters = input.nSliceRowClusters + rowStart;
679 auto*
buf = smem.getBuffer<V>(iWarp);
681 compressorMemcpyBuffered(
buf,
output->qTotU + rowsOffset, input.qTotU, nSectorRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0, compressor.
mMaxClusterFactorBase1024);
682 compressorMemcpyBuffered(
buf,
output->qMaxU + rowsOffset, input.qMaxU, nSectorRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0, compressor.
mMaxClusterFactorBase1024);
683 compressorMemcpyBuffered(
buf,
output->flagsU + rowsOffset, input.flagsU, nSectorRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0, compressor.
mMaxClusterFactorBase1024);
684 compressorMemcpyBuffered(
buf,
output->padDiffU + rowsOffset, input.padDiffU, nSectorRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0, compressor.
mMaxClusterFactorBase1024);
685 compressorMemcpyBuffered(
buf,
output->timeDiffU + rowsOffset, input.timeDiffU, nSectorRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0, compressor.
mMaxClusterFactorBase1024);
686 compressorMemcpyBuffered(
buf,
output->sigmaPadU + rowsOffset, input.sigmaPadU, nSectorRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0, compressor.
mMaxClusterFactorBase1024);
687 compressorMemcpyBuffered(
buf,
output->sigmaTimeU + rowsOffset, input.sigmaTimeU, nSectorRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0, compressor.
mMaxClusterFactorBase1024);
689 const uint16_t* nTrackClustersPtr = input.nTrackClusters + trackStart;
692 compressorMemcpyBuffered(
buf,
output->qTotA + tracksOffset, input.qTotA, nTrackClustersPtr, aClsFstIdx, tracksPerWarp, nLanes, iLane, 0);
693 compressorMemcpyBuffered(
buf,
output->qMaxA + tracksOffset, input.qMaxA, nTrackClustersPtr, aClsFstIdx, tracksPerWarp, nLanes, iLane, 0);
694 compressorMemcpyBuffered(
buf,
output->flagsA + tracksOffset, input.flagsA, nTrackClustersPtr, aClsFstIdx, tracksPerWarp, nLanes, iLane, 0);
695 compressorMemcpyBuffered(
buf,
output->sigmaPadA + tracksOffset, input.sigmaPadA, nTrackClustersPtr, aClsFstIdx, tracksPerWarp, nLanes, iLane, 0);
696 compressorMemcpyBuffered(
buf,
output->sigmaTimeA + tracksOffset, input.sigmaTimeA, nTrackClustersPtr, aClsFstIdx, tracksPerWarp, nLanes, iLane, 0);
699 uint32_t tracksOffsetDiff = tracksOffset - trackStart;
700 compressorMemcpyBuffered(
buf,
output->rowDiffA + tracksOffsetDiff, input.rowDiffA, nTrackClustersPtr, aClsFstIdx, tracksPerWarp, nLanes, iLane, 1);
701 compressorMemcpyBuffered(
buf,
output->sliceLegDiffA + tracksOffsetDiff, input.sliceLegDiffA, nTrackClustersPtr, aClsFstIdx, tracksPerWarp, nLanes, iLane, 1);
702 compressorMemcpyBuffered(
buf,
output->padResA + tracksOffsetDiff, input.padResA, nTrackClustersPtr, aClsFstIdx, tracksPerWarp, nLanes, iLane, 1);
703 compressorMemcpyBuffered(
buf,
output->timeResA + tracksOffsetDiff, input.timeResA, nTrackClustersPtr, aClsFstIdx, tracksPerWarp, nLanes, iLane, 1);
710 const auto& input = compressor.mPtrs;
711 auto*
output = compressor.mOutput;
713 const int32_t nWarps = nThreads / GPUCA_WARP_SIZE;
714 const int32_t iWarp = iThread / GPUCA_WARP_SIZE;
715 const int32_t nLanes = GPUCA_WARP_SIZE;
716 const int32_t iLane = iThread % GPUCA_WARP_SIZE;
720 compressorMemcpyBasic(
output->nSliceRowClusters, input.nSliceRowClusters, compressor.NSECTORS *
GPUCA_ROW_COUNT, nThreads, iThread);
721 compressorMemcpyBasic(
output->nTrackClusters, input.nTrackClusters, compressor.mMemory->nStoredTracks, nThreads, iThread);
722 compressorMemcpyBasic(
output->qPtA, input.qPtA, compressor.mMemory->nStoredTracks, nThreads, iThread);
723 compressorMemcpyBasic(
output->rowA, input.rowA, compressor.mMemory->nStoredTracks, nThreads, iThread);
724 compressorMemcpyBasic(
output->sliceA, input.sliceA, compressor.mMemory->nStoredTracks, nThreads, iThread);
725 compressorMemcpyBasic(
output->timeA, input.timeA, compressor.mMemory->nStoredTracks, nThreads, iThread);
726 compressorMemcpyBasic(
output->padA, input.padA, compressor.mMemory->nStoredTracks, nThreads, iThread);
727 }
else if (iBlock & 1) {
728 const uint32_t nGlobalWarps = nWarps * (nBlocks - 1) / 2;
729 const uint32_t iGlobalWarp = nWarps * (iBlock - 1) / 2 + iWarp;
732 uint32_t rowsPerWarp = (
nRows + nGlobalWarps - 1) / nGlobalWarps;
733 uint32_t rowStart = rowsPerWarp * iGlobalWarp;
734 uint32_t rowEnd = CAMath::Min(nRows, rowStart + rowsPerWarp);
735 if (rowStart >= nRows) {
739 rowsPerWarp = rowEnd - rowStart;
741 const uint32_t rowsOffset = calculateWarpOffsets(smem, input.nSliceRowClusters, rowStart, rowEnd, nWarps, iWarp, nLanes, iLane);
742 const uint32_t* clusterOffsets = &
clusters->clusterOffset[0][0] + rowStart;
743 const uint32_t* nSectorRowClusters = input.nSliceRowClusters + rowStart;
745 compressorMemcpyBuffered(
buf,
output->qTotU + rowsOffset, input.qTotU, nSectorRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0, compressor.mMaxClusterFactorBase1024);
746 compressorMemcpyBuffered(
buf,
output->qMaxU + rowsOffset, input.qMaxU, nSectorRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0, compressor.mMaxClusterFactorBase1024);
747 compressorMemcpyBuffered(
buf,
output->flagsU + rowsOffset, input.flagsU, nSectorRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0, compressor.mMaxClusterFactorBase1024);
748 compressorMemcpyBuffered(
buf,
output->padDiffU + rowsOffset, input.padDiffU, nSectorRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0, compressor.mMaxClusterFactorBase1024);
749 compressorMemcpyBuffered(
buf,
output->timeDiffU + rowsOffset, input.timeDiffU, nSectorRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0, compressor.mMaxClusterFactorBase1024);
750 compressorMemcpyBuffered(
buf,
output->sigmaPadU + rowsOffset, input.sigmaPadU, nSectorRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0, compressor.mMaxClusterFactorBase1024);
751 compressorMemcpyBuffered(
buf,
output->sigmaTimeU + rowsOffset, input.sigmaTimeU, nSectorRowClusters, clusterOffsets, rowsPerWarp, nLanes, iLane, 0, compressor.mMaxClusterFactorBase1024);
753 const uint32_t nGlobalWarps = nWarps * (nBlocks - 1) / 2;
754 const uint32_t iGlobalWarp = nWarps * (iBlock / 2 - 1) + iWarp;
756 const uint32_t nStoredTracks = compressor.mMemory->nStoredTracks;
757 uint32_t tracksPerWarp = (nStoredTracks + nGlobalWarps - 1) / nGlobalWarps;
758 uint32_t trackStart = tracksPerWarp * iGlobalWarp;
759 uint32_t trackEnd = CAMath::Min(nStoredTracks, trackStart + tracksPerWarp);
760 if (trackStart >= nStoredTracks) {
764 tracksPerWarp = trackEnd - trackStart;
766 const uint32_t tracksOffset = calculateWarpOffsets(smem, input.nTrackClusters, trackStart, trackEnd, nWarps, iWarp, nLanes, iLane);
767 const uint16_t* nTrackClustersPtr = input.nTrackClusters + trackStart;
768 const uint32_t* aClsFstIdx = compressor.mAttachedClusterFirstIndex + trackStart;
770 compressorMemcpyBuffered(
buf,
output->qTotA + tracksOffset, input.qTotA, nTrackClustersPtr, aClsFstIdx, tracksPerWarp, nLanes, iLane, 0);
771 compressorMemcpyBuffered(
buf,
output->qMaxA + tracksOffset, input.qMaxA, nTrackClustersPtr, aClsFstIdx, tracksPerWarp, nLanes, iLane, 0);
772 compressorMemcpyBuffered(
buf,
output->flagsA + tracksOffset, input.flagsA, nTrackClustersPtr, aClsFstIdx, tracksPerWarp, nLanes, iLane, 0);
773 compressorMemcpyBuffered(
buf,
output->sigmaPadA + tracksOffset, input.sigmaPadA, nTrackClustersPtr, aClsFstIdx, tracksPerWarp, nLanes, iLane, 0);
774 compressorMemcpyBuffered(
buf,
output->sigmaTimeA + tracksOffset, input.sigmaTimeA, nTrackClustersPtr, aClsFstIdx, tracksPerWarp, nLanes, iLane, 0);
777 uint32_t tracksOffsetDiff = tracksOffset - trackStart;
778 compressorMemcpyBuffered(
buf,
output->rowDiffA + tracksOffsetDiff, input.rowDiffA, nTrackClustersPtr, aClsFstIdx, tracksPerWarp, nLanes, iLane, 1);
779 compressorMemcpyBuffered(
buf,
output->sliceLegDiffA + tracksOffsetDiff, input.sliceLegDiffA, nTrackClustersPtr, aClsFstIdx, tracksPerWarp, nLanes, iLane, 1);
780 compressorMemcpyBuffered(
buf,
output->padResA + tracksOffsetDiff, input.padResA, nTrackClustersPtr, aClsFstIdx, tracksPerWarp, nLanes, iLane, 1);
781 compressorMemcpyBuffered(
buf,
output->timeResA + tracksOffsetDiff, input.timeResA, nTrackClustersPtr, aClsFstIdx, tracksPerWarp, nLanes, iLane, 1);
788 gatherBuffered<Vec32>(nBlocks, nThreads, iBlock, iThread, smem, processors);
794 gatherBuffered<Vec64>(nBlocks, nThreads, iBlock, iThread, smem, processors);
800 gatherBuffered<Vec128>(nBlocks, nThreads, iBlock, iThread, smem, processors);
806 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
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++)