Weird test failures

[chillenzer]

On the system at hand (HAL), the tests are all passing individually but fail if and only if "2D AccGpuCudaRt" and "3D AccGpuCudaRt" are both run in the same test run, so the following works

./tests <any test name>
./tests -f list.txt

whenever list.txt doesn't contain both of the above while the latter fails with any other list.txt, in particular

2D AccGpuCudaRt
3D AccGpuCudaRt

as a MWE. From what I can tell, neither the ordering of the two nor other tests in between change this observation. The error message from the ./tests is always

FAILED:
due to unexpected exception with message:
  /home/lenz93/workspace/src/mallocMC-makeReservationPolicyOwning/alpaka/
  include/alpaka/platform/PlatformUniformCudaHipRt.hpp(125) 'rc' A previous API
  call (not this one) set the error  : 'cudaErrorMisalignedAddress':
  'misaligned address'!

while

$ compute-sanitizer --show-backtrace=yes  --tool memcheck ./tests
========= COMPUTE-SANITIZER
========= Invalid __global__ atomic of size 4 bytes
=========     at atomicSub(unsigned int *, unsigned int)+0x4c0 in /opt/spack-modules/opt/spack/linux-ubuntu22.04-zen2/gcc-12.2.0/cuda-12.3.0-vhkdpnlaplsol7vzlt4wj55im3uemjkq/include/device_atomic_functions.hpp:122
=========     by thread (0,1,0) in block (0,3,1)
=========     Address 0x7ff15fe22051 is misaligned
=========     and is inside the nearest allocation at 0x7ff15fe00000 of size 1,048,576 bytes
=========     Device Frame:alpakaGlobal::AlpakaBuiltInAtomic<alpaka::AtomicSub, unsigned int, alpaka::hierarchy::Grids, void>::atomic(unsigned int *, unsigned int)+0x480 in /home/lenz93/workspace/src/mallocMC-makeReservationPolicyOwning/alpaka/include/alpaka/atomic/AtomicUniformCudaHip.hpp:161
[...]

(sometimes also out of bounds).

Gonna investigate further in the near future but any comments and hints are welcome!

[psychocoderHPC]

Even if you execute twice the 2D or 3D case it will fail. The reason is that the distribution policy XMallocSIMD has a bug.

Chnaging

mallocMC/tests/dimensions.cpp

Line 146 in bffe2aa

mallocMC::DistributionPolicies::XMallocSIMD<DistributionConfig>,

to mallocMC::DistributionPolicies::Noop, will fix the test.

[psychocoderHPC]

The problem in the XMallocSIMD policy is how we communicate between threads in a warp.

mallocMC/src/include/mallocMC/distributionPolicies/XMallocSIMD.hpp

Line 135 in bffe2aa

threadcount = popc(ballot(coalescible));

Since Volta threads in a warp can diverge, we should use in warp communication instead of shared memory.
One problem is that alpakas in warp thread shuffel is a collective warp operation. So we need to use native CUDA/ROCm calls.
An example can be found in PIConGPUs aggregated atomic increment

https://github.com/ComputationalRadiationPhysics/picongpu/blob/7edbba8d21172f97cad0e6bd84fdb93098773844/include/pmacc/kernel/atomic.hpp#L199-L204

[chillenzer]

Great catch! Thanks! Gonna try and fix it but not sure if I'll make it before the holidays.