TPC GPU Decoding code duplication removal

[cima22]

Created new class TPCClusterDecompressionCore.inc which contains common logic for old and new TPC track-model decoding. Same performance as before, less code.

[github-actions]

REQUEST FOR PRODUCTION RELEASES:
To request your PR to be included in production software, please add the corresponding labels called "async-" to your PR. Add the labels directly (if you have the permissions) or add a comment of the form (note that labels are separated by a ",")

+async-label <label1>, <label2>, !<label3> ...

This will add <label1> and <label2> and removes <label3>.

The following labels are available
async-2023-pbpb-apass4
async-2023-pp-apass4
async-2024-pp-apass1
async-2022-pp-apass7
async-2024-pp-cpass0
async-2024-PbPb-cpass0
async-2024-PbPb-apass1
async-2024-ppRef-apass1

[alibuild]

Error while checking build/O2/fullCI for 77d3702 at 2024-11-15 15:37:

## sw/BUILD/O2-latest/log
c++: error: unrecognized command-line option '--rtlib=compiler-rt'
c++: error: unrecognized command-line option '--rtlib=compiler-rt'
/sw/SOURCES/O2/13696-slc8_x86-64/0/GPU/GPUTracking/DataCompression/TPCClusterDecompressionCore.inc:158:130: error: no template named 'function' in namespace 'std'
/sw/SOURCES/O2/13696-slc8_x86-64/0/GPU/GPUTracking/DataCompression/TPCClusterDecompressionCore.inc(105): error: calling a __device__ function("o2::gpu::GPUTPCCompressionTrackModel::Mirror()") from a __host__ __device__ function("decompressTrack") is not allowed
/sw/SOURCES/O2/13696-slc8_x86-64/0/GPU/GPUTracking/DataCompression/TPCClusterDecompressionCore.inc(108): error: calling a __device__ function("_ZNK2o23gpu14GPUTPCGeometry5Row2XE1?") from a __host__ __device__ function("decompressTrack") is not allowed
/sw/SOURCES/O2/13696-slc8_x86-64/0/GPU/GPUTracking/DataCompression/TPCClusterDecompressionCore.inc(108): error: calling a __device__ function("_ZN2o23gpu27GPUTPCCompressionTrackModel9PropagateE1?1?") from a __host__ __device__ function("decompressTrack") is not allowed
/sw/SOURCES/O2/13696-slc8_x86-64/0/GPU/GPUTracking/DataCompression/TPCClusterDecompressionCore.inc(115): error: calling a __device__ function("o2::gpu::GPUTPCCompressionTrackModel::Z() const") from a __host__ __device__ function("decompressTrack") is not allowed
/sw/SOURCES/O2/13696-slc8_x86-64/0/GPU/GPUTracking/DataCompression/TPCClusterDecompressionCore.inc(115): error: calling a __device__ function("_ZN2o23gpu14GPUTPCGeometry12LinearZ2TimeE1?1?") from a __host__ __device__ function("decompressTrack") is not allowed
/sw/SOURCES/O2/13696-slc8_x86-64/0/GPU/GPUTracking/DataCompression/TPCClusterDecompressionCore.inc(115): error: calling a __device__ function("_ZN2o23tpc13ClusterNative8packTimeE1?") from a __host__ __device__ function("decompressTrack") is not allowed
/sw/SOURCES/O2/13696-slc8_x86-64/0/GPU/GPUTracking/DataCompression/TPCClusterDecompressionCore.inc(116): error: calling a __device__ function("_ZNK2o23gpu14GPUTPCGeometry5NPadsE1?") from a __host__ __device__ function("decompressTrack") is not allowed
/sw/SOURCES/O2/13696-slc8_x86-64/0/GPU/GPUTracking/DataCompression/TPCClusterDecompressionCore.inc(116): error: calling a __device__ function("o2::gpu::GPUTPCCompressionTrackModel::Y() const") from a __host__ __device__ function("decompressTrack") is not allowed
/sw/SOURCES/O2/13696-slc8_x86-64/0/GPU/GPUTracking/DataCompression/TPCClusterDecompressionCore.inc(116): error: calling a __device__ function("_ZNK2o23gpu14GPUTPCGeometry11LinearY2PadE1?1?1?") from a __host__ __device__ function("decompressTrack") is not allowed
/sw/SOURCES/O2/13696-slc8_x86-64/0/GPU/GPUTracking/DataCompression/TPCClusterDecompressionCore.inc(117): error: calling a __device__ function("_ZN2o23tpc13ClusterNative7packPadE1?") from a __host__ __device__ function("decompressTrack") is not allowed
/sw/SOURCES/O2/13696-slc8_x86-64/0/GPU/GPUTracking/DataCompression/TPCClusterDecompressionCore.inc(120): error: calling a __device__ function("_ZNK2o23gpu14GPUTPCGeometry5NPadsE1?") from a __host__ __device__ function("decompressTrack") is not allowed
/sw/SOURCES/O2/13696-slc8_x86-64/0/GPU/GPUTracking/DataCompression/TPCClusterDecompressionCore.inc(124): error: calling a __device__ function("_ZNK2o23gpu14GPUTPCGeometry5NPadsE1?") from a __host__ __device__ function("decompressTrack") is not allowed
/sw/SOURCES/O2/13696-slc8_x86-64/0/GPU/GPUTracking/DataCompression/TPCClusterDecompressionCore.inc(138): error: calling a __device__ function("_ZN2o23gpu27TPCClusterDecompressionCore20decompressTrackStoreE1?1?1?1?1?1?1?") from a __host__ __device__ function("decompressTrack") is not allowed
/sw/SOURCES/O2/13696-slc8_x86-64/0/GPU/GPUTracking/DataCompression/TPCClusterDecompressionCore.inc(139): error: calling a __device__ function("o2::tpc::ClusterNative::getPad() const") from a __host__ __device__ function("decompressTrack") is not allowed
/sw/SOURCES/O2/13696-slc8_x86-64/0/GPU/GPUTracking/DataCompression/TPCClusterDecompressionCore.inc(139): error: calling a __device__ function("_ZNK2o23gpu14GPUTPCGeometry11LinearPad2YE1?1?1?") from a __host__ __device__ function("decompressTrack") is not allowed
/sw/SOURCES/O2/13696-slc8_x86-64/0/GPU/GPUTracking/DataCompression/TPCClusterDecompressionCore.inc(140): error: calling a __device__ function("o2::tpc::ClusterNative::getTime() const") from a __host__ __device__ function("decompressTrack") is not allowed
/sw/SOURCES/O2/13696-slc8_x86-64/0/GPU/GPUTracking/DataCompression/TPCClusterDecompressionCore.inc(140): error: calling a __device__ function("_ZN2o23gpu14GPUTPCGeometry12LinearTime2ZE1?1?") from a __host__ __device__ function("decompressTrack") is not allowed
/sw/SOURCES/O2/13696-slc8_x86-64/0/GPU/GPUTracking/DataCompression/TPCClusterDecompressionCore.inc(143): error: calling a __device__ function("_ZNK2o23gpu14GPUTPCGeometry5Row2XE1?") from a __host__ __device__ function("decompressTrack") is not allowed
/sw/SOURCES/O2/13696-slc8_x86-64/0/GPU/GPUTracking/DataCompression/TPCClusterDecompressionCore.inc(143): error: calling a __device__ function("_ZN2o23gpu27GPUTPCCompressionTrackModel4InitE1?1?1?1?1?1?") from a __host__ __device__ function("decompressTrack") is not allowed
/sw/SOURCES/O2/13696-slc8_x86-64/0/GPU/GPUTracking/DataCompression/TPCClusterDecompressionCore.inc(145): error: calling a __device__ function("_ZN2o23gpu27GPUTPCCompressionTrackModel6FilterE1?1?1?") from a __host__ __device__ function("decompressTrack") is not allowed
/sw/SOURCES/O2/13696-slc8_x86-64/0/GPU/GPUTracking/DataCompression/TPCClusterDecompressionCore.inc(155): error: calling a __device__ function("_ZN2o23tpc13ClusterNativeC1E1?1?1?1?1?1?1?") from a __host__ __device__ function("decompressHitsStore") is not allowed
/sw/SOURCES/O2/13696-slc8_x86-64/0/GPU/GPUTracking/DataCompression/TPCClusterDecompressionCore.inc(160): error: calling a __device__ function("_ZN2o23tpc13ClusterNativeC1E1?1?1?1?1?1?1?") from a __host__ __device__ function("decompressHitsStore") is not allowed
/sw/SOURCES/O2/13696-slc8_x86-64/0/GPU/GPUTracking/DataCompression/TPCClusterDecompressionCore.inc(160): error: calling a __device__ function("_ZN2o23tpc13ClusterNativeC1E1?1?1?1?1?1?1?") from a __host__ __device__ function("decompressHitsStore") is not allowed
ninja: build stopped: subcommand failed.

Full log here.

[davidrohr]

Looks good, FullCI has passed now. You saw you validated that there is no performance regression neither for CPU nor for GPU? If you confirm, I'll merge it.