Validate NVHPC 21.9 compiler

[olupton]

We should make sure the next version of the NVIDIA HPC-SDK, 21.9, works and gives correct results.

Once this is done, we can update the deployed version on BB5, where we currently use 21.2. The three intermediate versions 21.3, 21.5 and 21.7 have bugs affecting CPU execution of GPU builds:

• 21.3 and 21.5 are affected by: Incorrect CPU results with #pragma acc atomic capture
• 21.7 is affected by: Miscompilation of simple CPU code with nvc/21.7

Another issue (Incorrect GPU results with #pragma acc atomic capture) appears to affect 21.2 and 21.7 (intermediate versions were not checked), but we seem to have a functioning workaround in BlueBrain/mod2c#68 and BlueBrain/nmodl#720.

Outstanding tasks:

• Fix compilation with 21.7: Tweak code to avoid NVHPC 21.7 ICE neuronsimulator/nrn#1385
• Wait for 21.9 release
  • fix any compilation issues
• Check if Avoid use of assert in OpenACC kernel along with use of atomic capture mod2c#68 and Avoid use of assert in the OpenACC kernel nmodl#720 are still needed.
• Check correctness of results, including non-CI tests like channel-benchmark

[olupton]

Unfortunately running the channel-benchmark tests with NVHPC 21.7 did not go well. This version of the nvc compiler appears to miscompile C code generated by NEURON's internal .mod file translator from the netstim_inhpoisson mechanism.

A minimal reproducer was sent to NVIDIA here: https://forums.developer.nvidia.com/t/miscompilation-of-simple-cpu-code-with-nvc-21-7/187517, we are waiting for their response.

[olupton]

Apparently the issue is fixed in the pre-release 21.9 compiler. I will rename this issue...

[olupton]

The 21.9 compiler is released (https://developer.nvidia.com/nvidia-hpc-sdk-219-downloads), so this issue can now move forwards.

[alkino]

it compiles and run basic test with 21.9

[olupton]

I opened BlueBrain/spack#1334 to install NVHPC 21.9 on BB5 and ran some tests using the modules from the PR CI.

After some confusion (it seems that having the llvm and nvhpc modules loaded at the same time is "more dangerous" with 21.9 than it was with 21.2, for reasons that I did not fully investigate but which are related to which omp.h is picked up), NEURON + CoreNEURON + NMODL + GPU built with with 21.9, and the full CTest suite passed (with the exception of needing a reference update to https://github.com/neuronsimulator/nrntest). This now includes the channel-benchmark tests.

[olupton]

To address the question about whether the "assert" bug is still there in 21.9, I reverted the merge commits of the workarounds we applied BlueBrain/mod2c#68 and BlueBrain/nmodl#720 and built with 21.2 and 21.9.

21.2 + NMODL:

coreneuron_modtests::test_watchrange_py_gpu (Failed)

which is the "expected" failure, see here for example.

21.9 + NMODL:

coreneuron_modtests::test_watchrange_py_gpu (Failed)
external_nrntest (Failed)

so it seems that the reported issue is still there with 21.9. The external_nrntest failure appears to be due to small numerical differences between 21.2 and 21.9 and can be fixed with a reference update (neuronsimulator/nrntest#21).

21.2 + MOD2C:

coreneuron_modtests::test_watchrange_py_gpu (Failed)

21.9 + MOD2C:

coreneuron_modtests::test_watchrange_py_gpu (Failed)

(neuronsimulator/nrntest#21 was included, so it is expected that external_nrntest did not fail).

Summary

These tests have not shown any new issues with NVHPC 21.9 (🎉), but the workarounds of BlueBrain/mod2c#68 and BlueBrain/nmodl#720 still seem to be required. I will update the thread on the NVIDIA forums (update: done).

Maybe @iomaganaris can check the olfactory bulb model with 21.9, and if that also looks OK we can move forward with BlueBrain/spack#1334 and bumping the default version on BB5 (and, afterwards, the CUDA version). It should be possible to get it setup on BB5 with:

#!/bin/bash
export MODULEPATH=/gpfs/bbp.cscs.ch/ssd/apps/hpc/jenkins/modules/all
module purge
# Cannot leave the LLVM module loaded because it messes with include paths and breaks NVHPC.
CLANG_FORMAT_PATH=$(module load unstable llvm/12.0.0 && dirname $(command -v clang-format))
module use /gpfs/bbp.cscs.ch/ssd/apps/hpc/jenkins/pulls/1334/modules/all
module load unstable gcc/9.3.0 nvhpc/21.9 cuda/11.0.2 hpe-mpi boost python-dev git flex bison ninja ccache cmake
export PATH=${CLANG_FORMAT_PATH}:${PATH}

[iomaganaris]

Everything got installed and worked fine for olfactory-bulb-3d model. With mod2c there are only some very small time shifts (1-2 timesteps) in ~40 spikes after ~800ms of simulation between nvhpc/21.2 and nvhpc/21.9. With nmodl+sympy there are time shifts in only 2 spikes after 1000ms. I believe that it is a very long sim and maybe they changed something related to the computation order/instructions so all we see is floating point differences.

[olupton]

I also tried NVHPC 21.9 with CUDA 11.4.2 and with BlueBrain/nmodl#720 reverted. All the tests passed, so it seems that the newer CUDA version does fix the issue.

[olupton]

With BlueBrain/spack#1334 we deployed 21.9 on BB5, and #679 fixed a small issue that the CI missed. I think we can close this.