Skip to content

Use cuComplex instead of thrust::complex #6

Closed

Description

@roiser

Done (SR)

available in

eemumu_auto/mg5Complex

Need to include also the complex implementation by Laurence

Activity

  1. added
    enhancementA feature we want to develop
    upstreamReady to be included in the MG5 code generator
    on Aug 12, 2020
  2. self-assigned this
    on Aug 13, 2020
  3. valassi commented on Aug 13, 2020

    @valassi
    Member

    This is now fully integrated here: roiser@bf1b625

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

  4. valassi commented on Aug 14, 2020

    @valassi
    Member

    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

  5. valassi commented on Aug 19, 2020

    @valassi
    Member

    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.
    image

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

    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...)
    image

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

    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
  6. valassi commented on Oct 21, 2021

    @valassi
    Member

    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

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Metadata

Metadata

Assignees

Labels

enhancementA feature we want to developupstreamReady to be included in the MG5 code generator

Type

No type

Projects

No projects

    Milestone

    No milestone

    Relationships

    None yet

    Development

    No branches or pull requests

    Issue actions