Leave single implementation of TRD RecoParam

[shahor02]

No description provided.

[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-apass1
async-2024-ppRef-apass1
async-2024-PbPb-apass2
async-2023-PbPb-apass5

[alibuild]

Error while checking build/O2/fullCI_slc9 for bbe74a0 at 2026-01-29 03:23:

## sw/BUILD/O2-latest/log
/sw/SOURCES/O2/15002-slc9_x86-64/0/GPU/GPUTracking/DataTypes/GPUTRDRecoParam.h(45): error: calling a __device__ function("o2::gpu::GPUTRDRecoParam::recalcTrkltCov(float, float, float, float *) const") from a __host__ function("recalcTrkltCov") is not allowed
/sw/SOURCES/O2/15002-slc9_x86-64/0/GPU/GPUTracking/DataTypes/GPUTRDRecoParam.h:45:5: error: no matching member function for call to 'recalcTrkltCov'
ninja: build stopped: subcommand failed.

Full log here.

[davidrohr]

looks principally good, but please move the recoParam from the constMemory to the calibObjects class, and some other minor comments.

[davidrohr]

This should go to GPUCalibObjectsConst calibObjects, not to GPUConstantMem directly.

[davidrohr]

You can use #include "GPUCommonArray.h" here, instad of conditional use of <array>

[shahor02]

@davidrohr ok, thanks. Could you explain why in the O2_fullCI these error are produced:

GPUTrackingCUDAExternalProvider.dir/GPUReconstructionCUDAExternalProvider.cu.o
/sw/SOURCES/O2/15002-slc9_x86-64/0/GPU/GPUTracking/DataTypes/GPUTRDRecoParam.h(45): error: calling a __device__ function("o2::gpu::GPUTRDRecoParam::recalcTrkltCov(float, float, float, float *) const") from a __host__ function("recalcTrkltCov") is not allowed
      recalcTrkltCov(tilt, snp, rowSize, cov.data());
      ^

1 error detected in the compilation of "/sw/SOURCES/O2/15002-slc9_x86-64/0/GPU/GPUTracking/Base/cuda/GPUReconstructionCUDAExternalProvider.cu".

and

In file included from /sw/SOURCES/O2/15002-slc9_x86-64/0/GPU/GPUTracking/TRDTracking/GPUTRDTracker.h:26:
/sw/SOURCES/O2/15002-slc9_x86-64/0/GPU/GPUTracking/DataTypes/GPUTRDRecoParam.h:45:5: error: no matching member function for call to 'recalcTrkltCov'
   45 |     recalcTrkltCov(tilt, snp, rowSize, cov.data());
      |     ^~~~~~~~~~~~~~
/sw/SOURCES/O2/15002-slc9_x86-64/0/GPU/GPUTracking/DataTypes/GPUTRDRecoParam.h:43:8: note: candidate function not viable: no known conversion from 'pointer' (aka 'float *') to 'std::array<float, 3> &' for 4th argument
   43 |   void recalcTrkltCov(const float tilt, const float snp, const float rowSize, std::array<float, 3>& cov) const
      |        ^                                                                      ~~~~~~~~~~~~~~~~~~~~~~~~~
/sw/SOURCES/O2/15002-slc9_x86-64/0/GPU/GPUTracking/DataTypes/GPUTRDRecoParam.h:48:15: note: candidate function not viable: call to __device__ function from __host__ function
   48 |   GPUd() void recalcTrkltCov(const float tilt, const float snp, const float rowSize, float* cov) const;
      |               ^

Naively I would expect that both methods https://github.com/AliceO2Group/AliceO2/pull/15002/changes/BASE..bbe74a01ad5f7beb21e1bd3f25f994a76ed3de2e#diff-68c42730194d5f1d35ec20df45bff05c28142482e5f0fa05ee7c3d8cae273154R41-R48 should be visible on the host and only the version with the pointer available on the GPU.

[davidrohr]

Well, the file is basically compiled 3 times:

1. For the CPU by gcc: All GPUd() and GPUh() goes away.
2. By hipcc / nvcc for the device: #if !defined(GPUCA_GPUCODE_DEVICE) hide the function, so it is ok
3. By hipcc / nvcc for the host (even if that is never used, but the compiler still sees the code): And here now you have the problem, that the void recalcTrkltCov(const float tilt, const float snp, const float rowSize, std::array<float, 3>& cov) const function calls GPUd() void recalcTrkltCov(const float tilt, const float snp, const float rowSize, float* cov) const; which gives you the error.

You can either hide it from the gpucode completely, by using #ifndef GPUCA_GPUCODE, or you have to declare it also as GPUd(), so even if not used, the compiler can parse it correctly when doing the host compilation pass.

[davidrohr]

Please forward-declare here, but do not include the file if not needed. It is only needed to know the class name for the pointer below, but not the class layout.