Implementing multivariate polynomial class - #7979

Just very minor things. Very nice! Especially the MultivariatePolynomial. Did you make some timing tests compared to the explicit formula?

Hi @wiechula,
Thanks for the comments, I have changed it accordingly.

When performing the interpolation on CPU tested with googlebenchmark the compile time version is around 20% slower than the hardcoded formula and the run time version is around 17 times slower than the hardcoded formula.
However when measuring the time for the tracking (GPUTPCGMMergerTrackFit) on an NVIDIA GPU the difference seems negligible for the compile time version. See this pdf (the y axis and y axis errors are just arbitrary to better distinguish the values.)

You can also look at the output from godbolt (left panel compile time evaluation and right panel hard coded formula) to see the differences.

@shahor02 @davidrohr , I guess the error is unrelated. Can this be merged?

Hello @davidrohr @sgorbuno,

now an error appeared during the slc8-GPU test (https://ali-ci.cern.ch/alice-build-logs/AliceO2Group/AliceO2/7979/a4302bcf068fcc27beee25c8402df65390b2aed1/build_O2_fullCI/pretty.html#errors) and I am not quite sure why it happens: /sw/SOURCES/O2/7979-slc8_x86-64/0/GPU/TPCFastTransformation/MultivariatePolynomial.h:92:74: error: cannot initialize a parameter of type 'const __private float *' with an lvalue of type '__generic o2::gpu::MultivariatePolynomial::DataTParams *const __generic' (aka '__generic float *const __generic') GPUd() float eval(const float x[/*Dim*/]) const { return this->evalPol(mParams, x); }

When I look to similar functions running on GPU e.g.:

By looking at the definition

It is needed in some cases when you are passing references/pointers to non-local variables to a function, which automatically deduces a template parameter from the type of that variable. The whole story is super complicated and related to OpenCL memory address space qualifiers.

basically, in

templace <class T> void foo(T& bla) {...}

int x;
boid bar()
{
  int y;
  foo(x);
  foo(y);
}

foo(x) yields T = int& while foo(y) yields T = __global__ int&, which are incompatible.
particularly, the latter will not work in case you have

templace <class T> void foo(T& bla) {
T z;
}

because it cannot instantiate __global__ int z; as local variable. This is prevented by adding GPUgeneric() before, which avoids that __global__ becomes part of T.

Thanks for the nice explanation. I hope it will fix the error.

thx, now it compiles. Have you tested on the EPN whether there are performance implications with the GPU?

thx, now it compiles. Have you tested on the EPN whether there are performance implications with the GPU?

I havent done a performance test on EPNs. @wiechula Could you run such a test on EPNs?

Please note, we have switched EPN access to using the slurm scheduler. Some instructions can be found here: https://indico.cern.ch/event/1111506/contributions/4675857/attachments/2370168/4047932/2022-01-01%20EPN%20access%20via%20Scheduler.pdf

I have added Jens to the users that may submit to the prod partition, so that he can use the GPUs.

Any update on the performance test of the new class on the EPN?

@davidrohr , here the comparison
Black: no dE/dx calculation
Red: dE/dx without any calibrations applied
Orange: dE/dx with Kr map and the new topological correction of this PR
The overhead of applying these calibrations during the dE/dx is ~4.5%, the overhead of dE/dx with calibrations in the tracking is ~30%.
gpuStats.pdf

@davidrohr , here the comparison Black: no dE/dx calculation Red: dE/dx without any calibrations applied Orange: dE/dx with Kr map and the new topological correction of this PR The overhead of applying these calibrations during the dE/dx is ~4.5%, the overhead of dE/dx with calibrations in the tracking is ~30%. gpuStats.pdf

Thanks @wiechula. Is it then ok to merge this @davidrohr ?