Skip to content

Commit 5eb4cda

Browse files
ChSonnabendChSonnabend
committed
Updating with separate kernel
1 parent 47d866b commit 5eb4cda

6 files changed

Lines changed: 210 additions & 4 deletions

File tree

‎GPU/GPUTracking/Definitions/Parameters/GPUParameters.csv‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -84,6 +84,7 @@ GPUTPCNNClusterizerKernels_determineClass2Labels,"""GPUCA_LB_GPUTPCNNClusterizer
8484
GPUTPCNNClusterizerKernels_publishClass1Regression,"""GPUCA_LB_GPUTPCNNClusterizerKernels""",,,,,,,,,,,,,,,
8585
GPUTPCNNClusterizerKernels_publishClass2Regression,"""GPUCA_LB_GPUTPCNNClusterizerKernels""",,,,,,,,,,,,,,,
8686
GPUTPCNNClusterizerKernels_publishDeconvolutionFlags,"""GPUCA_LB_GPUTPCNNClusterizerKernels""",,,,,,,,,,,,,,,
87+
GPUTPCNNClusterizerKernels_publishClass1RegressionWithNNDirection,"""GPUCA_LB_GPUTPCNNClusterizerKernels""",,,,,,,,,,,,,,,
8788
GPUTPCCFStreamCompaction_scanStart,"""GPUCA_PAR_CF_SCAN_WORKGROUP_SIZE""",,,,,,,,,,,,,,,"""GPUCA_PAR_CF_SCAN_WORKGROUP_SIZE"""
8889
GPUTPCCFStreamCompaction_scanUp,"""GPUCA_PAR_CF_SCAN_WORKGROUP_SIZE""",,,,,,,,,,,,,,,"""GPUCA_PAR_CF_SCAN_WORKGROUP_SIZE"""
8990
GPUTPCCFStreamCompaction_scanTop,"""GPUCA_PAR_CF_SCAN_WORKGROUP_SIZE""",,,,,,,,,,,,,,,"""GPUCA_PAR_CF_SCAN_WORKGROUP_SIZE"""

‎GPU/GPUTracking/Global/GPUChainTrackingClusterizer.cxx‎

Lines changed: 9 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1350,7 +1350,15 @@ int32_t GPUChainTracking::RunTPCClusterizer(bool synchronizeOutput)
13501350
runKernel<GPUTPCNNClusterizerKernels, GPUTPCNNClusterizerKernels::determineClass2Labels>({GetGrid(iSize, lane), krnlRunRangeNone}, iSector, clustererNNShadow.mNnInferenceOutputDType, propagateMCLabels, batchStart); // Assigning class labels
13511351
}
13521352
if (!clustererNNShadow.mNnClusterizerUseCfRegression) {
1353-
runKernel<GPUTPCNNClusterizerKernels, GPUTPCNNClusterizerKernels::publishClass1Regression>({GetGrid(iSize, lane), krnlRunRangeNone}, iSector, clustererNNShadow.mNnInferenceOutputDType, propagateMCLabels, batchStart); // Publishing class 1 regression results
1353+
if (buildNativeNNDirection) {
1354+
if (clustererShadow.mPclusterNNDirectionByRow != nullptr && clustererNNShadow.mNnClusterizerUseMomentumVector) {
1355+
runKernel<GPUTPCNNClusterizerKernels, GPUTPCNNClusterizerKernels::publishClass1RegressionWithNNDirection>({GetGrid(iSize, lane), krnlRunRangeNone}, iSector, clustererNNShadow.mNnInferenceOutputDType, propagateMCLabels, batchStart); // Publishing class 1 regression results with NN direction
1356+
} else {
1357+
runKernel<GPUTPCNNClusterizerKernels, GPUTPCNNClusterizerKernels::publishClass1Regression>({GetGrid(iSize, lane), krnlRunRangeNone}, iSector, clustererNNShadow.mNnInferenceOutputDType, propagateMCLabels, batchStart); // Publishing class 1 regression results
1358+
}
1359+
} else {
1360+
runKernel<GPUTPCNNClusterizerKernels, GPUTPCNNClusterizerKernels::publishClass1Regression>({GetGrid(iSize, lane), krnlRunRangeNone}, iSector, clustererNNShadow.mNnInferenceOutputDType, propagateMCLabels, batchStart); // Publishing class 1 regression results
1361+
}
13541362
if (nnApplication.mModelClass.getNumOutputNodes()[0][1] > 1 && nnApplication.mModelReg2.isInitialized()) {
13551363
runKernel<GPUTPCNNClusterizerKernels, GPUTPCNNClusterizerKernels::publishClass2Regression>({GetGrid(iSize, lane), krnlRunRangeNone}, iSector, clustererNNShadow.mNnInferenceOutputDType, propagateMCLabels, batchStart); // Publishing class 2 regression results
13561364
}

‎GPU/GPUTracking/TPCClusterFinder/GPUTPCNNClusterizerHost.cxx‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -138,7 +138,7 @@ void GPUTPCNNClusterizerHost::initClusterizer(const GPUSettingsProcessingNNclust
138138
if (!settings.nnClusterizerUseCfRegression) {
139139
if (mModelClass.getNumOutputNodes()[0][1] == 1 || !mModelReg2.isInitialized()) {
140140
clustererNN.mNnClusterizerModelReg1NumOutputNodes = mModelReg1.getNumOutputNodes()[0][1];
141-
if (clustererNN.mNnClusterizerModelReg1NumOutputNodes > 5) {
141+
if (clustererNN.mNnClusterizerModelReg1NumOutputNodes > 6) {
142142
clustererNN.mNnClusterizerUseMomentumVector = true;
143143
}
144144
} else {

‎GPU/GPUTracking/TPCClusterFinder/GPUTPCNNClusterizerKernels.cxx‎

Lines changed: 196 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -503,7 +503,202 @@ GPUdii() void GPUTPCNNClusterizerKernels::Thread<GPUTPCNNClusterizerKernels::pub
503503
} else if (clusterer.mPclusterPosInRow) {
504504
rowIndex = clusterer.mPclusterPosInRow[full_glo_idx];
505505
}
506-
if (clustererNN.mNnClusterizerUseMomentumVector && clusterer.mPclusterNNDirectionByRow != nullptr && rowIndex < clusterer.mNMaxClusterPerRow) {
506+
CPU_ONLY(labelAcc->commit(peak.row(), rowIndex, clusterer.mNMaxClusterPerRow));
507+
}
508+
509+
template <>
510+
GPUdii() void GPUTPCNNClusterizerKernels::Thread<GPUTPCNNClusterizerKernels::publishClass1RegressionWithNNDirection>(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUSharedMemory& smem, processorType& processors, uint8_t sector, int8_t dtype, int8_t withMC, uint32_t batchStart)
511+
{
512+
uint32_t glo_idx = get_global_id(0);
513+
auto& clusterer = processors.tpcClusterer[sector];
514+
auto& clustererNN = processors.tpcNNClusterer[sector];
515+
if (glo_idx >= (uint32_t)clustererNN.mNnClusterizerBatchedMode) {
516+
return;
517+
}
518+
519+
uint32_t maxClusterNum = clusterer.mPmemory->counters.nClusters;
520+
uint32_t full_glo_idx = glo_idx + batchStart;
521+
int32_t model_output_index = glo_idx * clustererNN.mNnClusterizerModelReg1NumOutputNodes;
522+
523+
CfArray2D<PackedCharge> chargeMap(reinterpret_cast<PackedCharge*>(clusterer.mPchargeMap));
524+
uint32_t peakIndex = CAMath::Min(full_glo_idx, maxClusterNum - 1);
525+
CfChargePos peak = clusterer.mPfilteredPeakPositions[peakIndex];
526+
float central_charge = static_cast<float>(chargeMap[peak].unpack());
527+
528+
CPU_ONLY(MCLabelAccumulator labelAccElem(clusterer));
529+
MCLabelAccumulator* labelAcc = CPU_PTR(&labelAccElem);
530+
531+
if (full_glo_idx >= maxClusterNum) {
532+
if (withMC) {
533+
ClusterAccumulator dummy_pc;
534+
CPU_ONLY(labelAcc->collect(peak, central_charge));
535+
GPUTPCCFClusterizer::buildCluster(
536+
clusterer.Param().rec,
537+
chargeMap,
538+
peak,
539+
smem.posBcast,
540+
smem.buf,
541+
smem.innerAboveThreshold,
542+
&dummy_pc,
543+
labelAcc);
544+
}
545+
return;
546+
}
547+
548+
tpc::ClusterNative* clusterOut = clusterer.mPclusterByRow;
549+
550+
ClusterAccumulator pc;
551+
552+
if (withMC) {
553+
ClusterAccumulator dummy_pc;
554+
CPU_ONLY(labelAcc->collect(peak, central_charge));
555+
GPUTPCCFClusterizer::buildCluster(
556+
clusterer.Param().rec,
557+
chargeMap,
558+
peak,
559+
smem.posBcast,
560+
smem.buf,
561+
smem.innerAboveThreshold,
562+
&dummy_pc,
563+
labelAcc);
564+
}
565+
if ((clusterer.mPmemory->fragment).isOverlap(peak.time())) {
566+
if (clusterer.mPclusterPosInRow) {
567+
clusterer.mPclusterPosInRow[full_glo_idx] = clusterer.mNMaxClusterPerRow;
568+
}
569+
return;
570+
}
571+
572+
bool notSinglePad = false, notSingleTime = false;
573+
for (uint16_t i = 0; i < 8; i++) {
574+
Delta2 d = cfconsts::InnerNeighbors[i];
575+
CfChargePos tmp_pos = peak.delta(d);
576+
float v = static_cast<float>(chargeMap[tmp_pos].unpack());
577+
notSinglePad |= (d.x != 0) && (v > 0.f);
578+
notSingleTime |= (d.y != 0) && (v > 0.f);
579+
}
580+
581+
float publishPadPosition = 0.f, publishTimePosition = 0.f;
582+
if (dtype == 0) {
583+
publishPadPosition = static_cast<float>(peak.pad()) + clustererNN.mOutputDataReg1_32[model_output_index];
584+
publishTimePosition = static_cast<float>(peak.time()) + clustererNN.mOutputDataReg1_32[model_output_index + 1];
585+
isBoundaryPublish(full_glo_idx, static_cast<int32_t>(peak.row()), publishPadPosition, publishTimePosition);
586+
pc.setFull(central_charge * clustererNN.mOutputDataReg1_32[model_output_index + 4],
587+
publishPadPosition,
588+
notSinglePad ? clustererNN.mOutputDataReg1_32[model_output_index + 2] : 0.f,
589+
(clusterer.mPmemory->fragment).start + publishTimePosition,
590+
notSingleTime ? clustererNN.mOutputDataReg1_32[model_output_index + 3] : 0.f,
591+
clustererNN.mClusterFlags[2 * glo_idx],
592+
clustererNN.mClusterFlags[2 * glo_idx + 1]);
593+
} else {
594+
publishPadPosition = static_cast<float>(peak.pad()) + clustererNN.mOutputDataReg1_16[model_output_index].ToFloat();
595+
publishTimePosition = static_cast<float>(peak.time()) + clustererNN.mOutputDataReg1_16[model_output_index + 1].ToFloat();
596+
isBoundaryPublish(full_glo_idx, static_cast<int32_t>(peak.row()), publishPadPosition, publishTimePosition);
597+
pc.setFull(central_charge * clustererNN.mOutputDataReg1_16[model_output_index + 4].ToFloat(),
598+
publishPadPosition,
599+
notSinglePad ? clustererNN.mOutputDataReg1_16[model_output_index + 2].ToFloat() : 0.f,
600+
(clusterer.mPmemory->fragment).start + publishTimePosition,
601+
notSingleTime ? clustererNN.mOutputDataReg1_16[model_output_index + 3].ToFloat() : 0.f,
602+
clustererNN.mClusterFlags[2 * glo_idx],
603+
clustererNN.mClusterFlags[2 * glo_idx + 1]);
604+
}
605+
606+
// if (boundaryFlag != 0) { // Prints the entire NN input for the given index
607+
// // Build a simple buffer manually (float with 3 decimals)
608+
// const int MAX_CHARS = 4096;
609+
// char buffer[MAX_CHARS];
610+
// int pos = 0;
611+
//
612+
// auto appendChar = [&](char c) {
613+
// if (pos < MAX_CHARS - 1) buffer[pos++] = c;
614+
// };
615+
// auto appendStr = [&](const char* s) {
616+
// while (*s && pos < MAX_CHARS - 1) buffer[pos++] = *s++;
617+
// };
618+
// auto appendUInt = [&](uint32_t v) {
619+
// char tmp[16]; int tp = 0;
620+
// if (v == 0) { appendChar('0'); return; }
621+
// while (v && tp < 16) { tmp[tp++] = char('0' + (v % 10)); v /= 10; }
622+
// while (tp--) appendChar(tmp[tp]);
623+
// };
624+
// auto appendInt = [&](int v) {
625+
// if (v < 0) { appendChar('-'); v = -v; }
626+
// appendUInt((uint32_t)v);
627+
// };
628+
// auto appendFloat = [&](float f) {
629+
// if (f < 0) { appendChar('-'); f = -f; }
630+
// int ip = (int)f;
631+
// float frac = f - (float)ip;
632+
// appendInt(ip);
633+
// appendChar('.');
634+
// for (int i = 0; i < 3; i++) {
635+
// frac *= 10.f;
636+
// int d = (int)frac;
637+
// appendChar((char)('0' + (d < 0 ? 0 : (d > 9 ? 9 : d))));
638+
// frac -= d;
639+
// if (frac < 0) frac = 0;
640+
// }
641+
// };
642+
//
643+
// appendStr("(NN CLUS) DEBUG: Boundary cluster detected (sector ");
644+
// appendUInt(sector);
645+
// appendStr(", row ");
646+
// appendUInt(peak.row());
647+
// appendStr(", pad ");
648+
// appendFloat(publishPadPosition);
649+
// appendStr(", time ");
650+
// appendFloat(publishTimePosition);
651+
// appendStr(") [glo_idx=");
652+
// appendUInt(glo_idx);
653+
// appendStr(" elemSize=");
654+
// appendInt(clustererNN.mNnClusterizerElementSize);
655+
// appendStr(" dtype=");
656+
// appendInt(dtype);
657+
// appendStr("] INPUT:");
658+
//
659+
// int elemSize = clustererNN.mNnClusterizerElementSize;
660+
// int baseIdx = glo_idx * elemSize;
661+
//
662+
// int maxPrint = elemSize;
663+
// for (int i = 0; i < maxPrint; ++i) {
664+
// appendChar(' ');
665+
// float v = (dtype == 0) ? clustererNN.mInputData_16[baseIdx + i].ToFloat()
666+
// : clustererNN.mInputData_32[baseIdx + i];
667+
// appendFloat(v);
668+
// if (pos > (MAX_CHARS - 32)) { appendStr(" ..."); break; }
669+
// }
670+
//
671+
// buffer[pos] = 0;
672+
// printf("%s\n", buffer);
673+
// }
674+
675+
tpc::ClusterNative myCluster;
676+
bool rejectCluster = !pc.toNative(peak, central_charge, myCluster, clusterer.Param(), chargeMap);
677+
if (clustererNN.mNnClusterizerUseClassification) {
678+
rejectCluster |= (clustererNN.mOutputDataClass[peakIndex] <= 0);
679+
}
680+
if (rejectCluster) {
681+
if (clusterer.mPclusterPosInRow) {
682+
clusterer.mPclusterPosInRow[full_glo_idx] = clusterer.mNMaxClusterPerRow;
683+
}
684+
return;
685+
}
686+
uint32_t rowIndex = 0;
687+
if (clusterOut != nullptr) {
688+
rowIndex = GPUTPCCFClusterizer::sortIntoBuckets(
689+
clusterer,
690+
myCluster,
691+
peak.row(),
692+
clusterer.mNMaxClusterPerRow,
693+
clusterer.mPclusterInRow,
694+
clusterOut);
695+
if (clusterer.mPclusterPosInRow != nullptr) {
696+
clusterer.mPclusterPosInRow[full_glo_idx] = rowIndex;
697+
}
698+
} else if (clusterer.mPclusterPosInRow) {
699+
rowIndex = clusterer.mPclusterPosInRow[full_glo_idx];
700+
}
701+
if (rowIndex < clusterer.mNMaxClusterPerRow) {
507702
clusterer.mPclusterNNDirectionByRow[peak.row() * clusterer.mNMaxClusterPerRow + rowIndex] = getClass1NNDirection(clustererNN, dtype, model_output_index);
508703
}
509704
CPU_ONLY(labelAcc->commit(peak.row(), rowIndex, clusterer.mNMaxClusterPerRow));

‎GPU/GPUTracking/TPCClusterFinder/GPUTPCNNClusterizerKernels.h‎

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -53,7 +53,8 @@ class GPUTPCNNClusterizerKernels : public GPUKernelTemplate
5353
determineClass2Labels = 4,
5454
publishClass1Regression = 5,
5555
publishClass2Regression = 6,
56-
publishDeconvolutionFlags = 7
56+
publishDeconvolutionFlags = 7,
57+
publishClass1RegressionWithNNDirection = 8
5758
};
5859

5960
template <int32_t iKernel = defaultKernel, typename... Args>

‎GPU/GPUTracking/kernels.cmake‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -135,6 +135,7 @@ o2_gpu_add_kernel("GPUTPCNNClusterizerKernels, determineClass2Labels" "= TPC
135135
o2_gpu_add_kernel("GPUTPCNNClusterizerKernels, publishClass1Regression" "= TPCNNCLUSTERFINDER" LB uint8_t sector int8_t dtype int8_t withMC uint32_t batchStart)
136136
o2_gpu_add_kernel("GPUTPCNNClusterizerKernels, publishClass2Regression" "= TPCNNCLUSTERFINDER" LB uint8_t sector int8_t dtype int8_t withMC uint32_t batchStart)
137137
o2_gpu_add_kernel("GPUTPCNNClusterizerKernels, publishDeconvolutionFlags" "= TPCNNCLUSTERFINDER" LB uint8_t sector int8_t dtype int8_t withMC uint32_t batchStart)
138+
o2_gpu_add_kernel("GPUTPCNNClusterizerKernels, publishClass1RegressionWithNNDirection" "= TPCNNCLUSTERFINDER" LB uint8_t sector int8_t dtype int8_t withMC uint32_t batchStart)
138139
endif()
139140

140141
o2_gpu_kernel_add_parameter(NEIGHBOURS_FINDER_MAX_NNEIGHUP # Number of neighhbours finder hits to cache in shared memory

0 commit comments

Comments
 (0)