Memory layout (shared/global/local, AOSOA...) for intermediate wavefunctions in ME calculations

[roiser]

eemumu_AV/master

AV should check for bugs in global memory implementation (seems not possible to exhaust it with eemumu…)

[valassi]

This is ongoing. I fixed the global memory issue in the meantime: roiser@649bdc4

But I still get much worse performance than with local memory or shared memory. It is even worse than before, because now I can use 16384 blocks and the performance penalty of global vs local is much worse at 16384 clocks than it was at 256.

I would say that this remains to be better undesrtoood and profiled

[valassi]

I have closed the AOS/SOA issue #16 fo rthe memory layout of momenta (and random numbers). Now the access to momenta in memory seems well coalesced... and this does not even seem to have a big impact.

I have renamed this issue #7 as "Memory layout (shared/global/local, AOSOA...) for intermediate wavefunctions in ME calculations". It may still be that optimising those memory structures like w[5][6] will have an impact.

In other words, the memory of the input data (momenta) to signmakin and of output data (MEs) seems now under control, and maybe the layout is not that relevant. We should rather focus on the memory layout inside the sigmakin calculation, ie nside the ixxx, oxxx and FV functions.

[valassi]

I repeat here what I wrote in issue #26 about maxregcount.

I think that I got this whole idea of local/global/shared wrong, because I had not understood two things

• one, that 'registers' are vector registers, and that warps effectively do "simd" n those vector registers (so even the 'local' w[5][6] is already vectorizxed), and there is no need to vectorize them manually as in c++ on a cpu
• two, that what i read about 'local storage is thread-local global storage' is also in part a misconception, because those stack variables like w[5][6] stay on the registers as long as possible, and only spill out to "local" memory (i.e. thread-local global memory) if there is no space left in the registers.

It seems much better therefore to concentrate on the 'local' coding with w[5][6] and drop the shared and global implementations, which are complex and have worse performance. For instance

• In issue Use maxrregcount to reduce register usage and improve throughput #26 I investigated the use of maxregcount. This does improve things, but needs careful fine tuning. Essentially, now with the latest AOSOA structure (issue AOS/SOA for input particle 4-momenta (and random numbers) #16) we are doing minimal usage of the memory bandwidth. Setting a limit on the number of registers forces the spillover from registries to "local" memory, which in any case is not an issue as we have plenty of unused bandwidth. This allows increasing occupancy by reducing the number of registries.
• In issue Compile-time const (and constexpr) vs constant-memory constants (layouts; physics parameters?) #23, I wonder if moving some physics constants from constant memory to hardcoded executables would also reduce the number of registries.
• We should investigate individual operations like sqrt (issue Investigate how expensive sqrt is for CUDA (registers) and C++ (vectorization etc) #15)
• Clearly using float instead of double would not only allow the use of FP32, but would also reduce the number of registers (issue Use single precision or even half precision #5)
• I wonder if using a better complex implementation (issue Use cuComplex instead of thrust::complex #6) with home-made wrapping of doubles would also decrease register usage?
• And we should also consider splitting kernels into smaller chunks, for instance using cuda graphs (issue cuda graphs #12)

As far as this work on local vs global vcs storage is concerned, I would say that

• The shared implementation, as it is now for w[5][6], has no future. It has a 35% throughpout penalty and it already uses 15k bytes of shared memory with 32 threads, when there are only 64k (or maybe 48k?). It is already impossible to use 128 threads. With more complex processes like ggttgg. it would simply explode. This is not a useful way for us to use shared memory.
• The global memory implementation has a huge hit, like 70%. As discussed in issue Use maxrregcount to reduce register usage and improve throughput #26, it seems much better to let the system handlk the spill of some memory to "local" (thread-local global) memory, rather than do it ourselves.
• The only reason why we may need to resurrect some of these global memory studies in the future is if we do split the sigmakin into smaller chunks. In that case we need to store intermediate results (like the w[5][6]...) into global memory (NOT shared memory) to move it across kernels. But it would need a much more profound reengineering of the kernels.

In conclusion, I woudl simply remove the present global and shared implementations. THis will clean up tyhe code.

Before doing that, I will still post a few graphs from profiling of shared and global below.

[valassi]

First, compare SHARED to "LOCAL".

It uses memory more, but SM usage decreases (unlike the case of maxregcount 128, where using more the memory allows a better usage of SM too).

The roofline is not very interesting. Actually the point moves to the right, not to the left (arithm,etic intensity increases). Actually FP64 decreases, while the LSU load store unit increases a lot

The memory plot shows that shared memory is used, which was not the case before. There are also wuit ea few bank conflicts (I am not sure if they are many or if that's an issue)

On the stall plot, it seems that the "useful" stalls like math pipe throttle decrease a lot

The mix of instructions changes a lot. But I would say that there a lot of new, not essential operations that are added. So we do more computations maybe, but not in useful things?

Finally, the occupancy plot. This is probably one of the most useful plots here. All this work was to reduce the number of registers, and we do reduce them a bit from 176 to 164. But, because we are using shared memory, this does NOT translate into a higher occupancy, actually the overall occupancy is lower. (I am not sure what is involved, but this is what the bottom plot syas)

So, in conclusion this usage of shared memory done this way is useless.

[valassi]

About global memory. This looks a bit like the plot for maxreg=64. We have increased memory usage so much, that SM decreases and we have become memory bound

And again the roofline says it all, we are now on the diagonal

Not surprisingly, memory usage has increased enormously

And the warp plots are interesting, though I do not understand all of it. We actually have slightly higher active warps per scheduler, but much fewer issued warps per scheduler. And again the only "useful" stall, the math pipe throttle, has gone down and we are memory bound.

As in the SHARED case, again the instruction mix is very different, and we do more instructions but probably they are useless ones.. (ie some we could avoid in the ME computation)

Fianlly, this time the reduced number of registries (166 vs 176) does increase occupancy a bit, but this is completely irrelevant now as we have become memory bound

[valassi]

Ok I have removed all code for SHARED and GLOBAL.

In case, this is where it was done: d1a5097

This can be closed now