Use cuComplex instead of thrust::complex

[roiser]

Done (SR)

available in

eemumu_auto/mg5Complex

Need to include also the complex implementation by Laurence

[valassi]

This is now fully integrated here: roiser@bf1b625

But it does seem slower than thurst. See https://docs.google.com/document/d/1g2xwJ2FsSlxHvSUdPZjCyFW7zhsblMQ4g8UHlrkWyVw/edit#

[valassi]

When integrating cucomplex, I looked at the header and thought it might have a better performance because it includes explicit align, which maybe thrust does not.

This is an interesting read however: https://stackoverflow.com/questions/46965768/optimizing-memory-access-for-complex-numbers. There is a comment from the thrust developer there, saying that he implemented a better alignment in thrust.

PS About this specific issue: I am no longer convinced alignment is an issue for our complex, as registers are vector registers anyway... it may be relevant only if we spill to global memory and/or we fragment compution eg with cuda graphs

[valassi]

Before I move to other things, f=I sump also some plots on cucomplex vs thrust. I confirm that I observe a small degradation of throughput with cucomplex, 6.5E8 instead of 6.8E8.

First, a small degradation is visible in SOL, the SM is a bit worse, while memory is more used. Note the registers are 184 instead opf 172, which may also be the main reason.

The warp state actually seems better, with more math pipe throttle?

Well, actually the issue may be that there are many more instructions executed. This may depend on the types themselves, or maybe on the headers with operator overloading (which I copied from Stefan, but I made misnor modifications, it may be worth seeing if I made some transcription mistakes...)

Registers per thread are higher, but this does not seem to have any negative impact on occupancy

Conclusion:

• the two main differences seem to be the higher number of registers and the higher number of instructions
• the number of registers seems to have no negative impact
• the higher number of instructions may well be the only culprit for the observed slowdown, it may be worth checking if I made a transcription mistake (or judgement error) while copying/modifying Stefan's headers
• the difference is any case mninimal, and for the moment it's worth staying with thrust

[valassi]

I would say that this can be closed.

We have cucomplex still in the headers of the the code, even if they are not the default.

We can always test it again if we like. But actually I tend to think that a custom cxtype_s with a simple pair of numbers (like cxtype_v is for RRRRIIII) would be more interesting, and allows some custom returns by reference.

I am closing this