Re-enable CI tests on hackathon branch #717

olupton · 2021-12-14T08:49:55Z

Re-enable various bits of CI and don't try and enable OpenMP target offload with MOD2C.

With this change we have one new GitLab CI build, giving a total of:

CPU-only with the Intel compiler and MOD2C
CPU-only with the Intel compiler and NMODL, no SymPy
GPU-enabled using OpenACC, using MOD2C and OpenMP host parallelism
GPU-enabled using OpenACC, using NMODL and SymPy, no OpenMP host parallelism
GPU-enabled using OpenMP target offload, using NMODL, no SymPy (*), and OpenMP host parallelism

(*) SymPy doesn't work because of OpenMP/Eigen issues: https://forums.developer.nvidia.com/t/enabling-openmp-offload-breaks-openacc-code/196643

The other main change concerning OpenMP support with the NVIDIA compilers is that we no longer enable OpenACC. During the hackathon, when we had not migrated the data transfer code to use OpenMP, we were enabling both OpenACC and OpenMP (-acc -mp=gpu) and relying on the compilers' interoperability. This PR drops the -acc in that case. Making this work required numerous small fixes, with a lot of overlap with the draft changes for LLVM and XLC offload support.

Needs Fix NVHPC + OpenMP ~ OpenACC compilation nmodl#784.

Use certain branches for the SimulationStack CI

CI_BRANCHES:NEURON_BRANCH=master,NMODL_BRANCH=hackathon_main,

bbpbuildbot · 2021-12-14T10:45:34Z

bbpbuildbot · 2021-12-14T15:13:26Z

bbpbuildbot · 2021-12-14T16:06:07Z

olupton · 2021-12-14T16:59:40Z

The CI failures seem to be somehow related to differences between the phase 1 and phase 2 GPU nodes on the cluster. I can only reproduce the failure locally on a phase 2 node, and only with 2 MPI ranks.

bbpbuildbot · 2021-12-15T09:12:19Z

bbpbuildbot · 2021-12-15T15:05:17Z

…oad.

bbpbuildbot · 2021-12-16T10:29:22Z

coreneuron/gpu/nrn_acc_manager.cpp

iomaganaris

LGTM in general. Some small questions here and there.
Another question I have whether the OpenMP + Unified memory has been tested either in the CI or locally

coreneuron/mechanism/mech/mod2c_core_thread.hpp

coreneuron/permute/cellorder.cpp

coreneuron/sim/scopmath/ssimplic_thread.cpp

olupton · 2021-12-17T10:40:43Z

Another question I have whether the OpenMP + Unified memory has been tested either in the CI or locally

I think that the Jenkins CI running on #713 will be trying to do something there, but I have not checked it carefully. I never tried this myself locally.

Co-authored-by: Ioannis Magkanaris <[email protected]>

iomaganaris · 2021-12-17T11:05:51Z

Another question I have whether the OpenMP + Unified memory has been tested either in the CI or locally

I think that the Jenkins CI running on #713 will be trying to do something there, but I have not checked it carefully. I never tried this myself locally.

I see. Something to check before merging to master then

pramodk

LGTM!
(Skimmed through changes except CI part and LGTM. When the PR will be created against master then we can discuss overall changes.)

pramodk · 2021-12-17T11:21:42Z

coreneuron/gpu/nrn_acc_manager.cpp

+    // It seems that with NVHPC 21.9 then only setting the default OpenMP device
+    // is not enough: there were errors on some nodes when not-the-0th GPU was
+    // used. These seemed to be related to the NMODL instance structs, which are
+    // allocated using cudaMallocManaged.


not sure if this is still an issue but curious if this is still an issue.

Still an issue in 21.11 you mean?

bbpbuildbot · 2021-12-17T12:24:19Z

Summary of changes: - Support OpenMP target offload when NMODL and GPU support are enabled. (#693, #704, #705, #707, #708, #716, #719) - Use sensible defaults for the --nwarp parameter, improving the performance of the Hines solver with --cell-permute=2 on GPU. (#700, #710, #718) - Use a Boost memory pool, if Boost is available, to reduce the number of independent CUDA unified memory allocations used for Random123 stream objects. This speeds up initialisation of models using Random123, and also makes it feasible to use NSight Compute on models using Random123 and for NSight Systems to profile initialisation. (#702, #703) - Use -cuda when compiling with NVHPC and OpenACC or OpenMP, as recommended on the NVIDIA forums. (#721) - Do not compile for compute capability 6.0 by default, as this is not supported by NVHPC with OpenMP target offload. - Add new GitLab CI tests so we test CoreNEURON + NMODL with both OpenACC and OpenMP. (#698, #717) - Add CUDA runtime header search path explicitly, so we don't rely on it being implicit in our NVHPC localrc. - Cleanup unused code. (#711) Co-authored-by: Pramod Kumbhar <[email protected]> Co-authored-by: Ioannis Magkanaris <[email protected]> Co-authored-by: Christos Kotsalos <[email protected]> Co-authored-by: Nicolas Cornu <[email protected]>

Summary of changes: - Support OpenMP target offload when NMODL and GPU support are enabled. (BlueBrain/CoreNeuron#693, BlueBrain/CoreNeuron#704, BlueBrain/CoreNeuron#705, BlueBrain/CoreNeuron#707, BlueBrain/CoreNeuron#708, BlueBrain/CoreNeuron#716, BlueBrain/CoreNeuron#719) - Use sensible defaults for the --nwarp parameter, improving the performance of the Hines solver with --cell-permute=2 on GPU. (BlueBrain/CoreNeuron#700, BlueBrain/CoreNeuron#710, BlueBrain/CoreNeuron#718) - Use a Boost memory pool, if Boost is available, to reduce the number of independent CUDA unified memory allocations used for Random123 stream objects. This speeds up initialisation of models using Random123, and also makes it feasible to use NSight Compute on models using Random123 and for NSight Systems to profile initialisation. (BlueBrain/CoreNeuron#702, BlueBrain/CoreNeuron#703) - Use -cuda when compiling with NVHPC and OpenACC or OpenMP, as recommended on the NVIDIA forums. (BlueBrain/CoreNeuron#721) - Do not compile for compute capability 6.0 by default, as this is not supported by NVHPC with OpenMP target offload. - Add new GitLab CI tests so we test CoreNEURON + NMODL with both OpenACC and OpenMP. (BlueBrain/CoreNeuron#698, BlueBrain/CoreNeuron#717) - Add CUDA runtime header search path explicitly, so we don't rely on it being implicit in our NVHPC localrc. - Cleanup unused code. (BlueBrain/CoreNeuron#711) Co-authored-by: Pramod Kumbhar <[email protected]> Co-authored-by: Ioannis Magkanaris <[email protected]> Co-authored-by: Christos Kotsalos <[email protected]> Co-authored-by: Nicolas Cornu <[email protected]> CoreNEURON Repo SHA: BlueBrain/CoreNeuron@423ae6c

olupton force-pushed the olupton/reenable-ci branch from 7f5253c to 8542e9c Compare December 14, 2021 09:46

olupton mentioned this pull request Dec 14, 2021

Fix NVHPC + OpenMP ~ OpenACC compilation BlueBrain/nmodl#784

Merged

olupton closed this Dec 15, 2021

olupton reopened this Dec 15, 2021

olupton requested a review from iomaganaris December 15, 2021 15:44

olupton added 15 commits December 16, 2021 09:37

Restore clang/cmake formatting checks.

2c6a65b

Prefer OpenACC with MOD2C. Disable OpenACC with NMODL and OpenMP offl…

f24d75e

…oad.

Re-enable GitLab CI from master.

c576713

apply/disable clang-format

8a6dc8a

Add NMODL + OpenACC test.

05b993c

Don't pass --gpu in non-GPU builds.

3942035

clang-format

12f019c

Pass --mpi to channel-benchmark test.

9ca7296

Fix typo.

938a983

Try and fix OpenACC-disabled NVHPC builds.

ddb01ff

Convert more #pragma acc to nrn_pragma_acc(...).

3eb9e96

More fixups.

2b7c481

Juggle openmp and sympy variants.

c954524

Call cudaSetDevice in OpenMP mode.

16f9111

Use nrn_pragma_omp for indentation.

f4c2eb6

olupton force-pushed the olupton/reenable-ci branch from 7f060ad to f4c2eb6 Compare December 16, 2021 09:00

olupton requested review from alkino and pramodk December 17, 2021 09:08

alkino reviewed Dec 17, 2021

View reviewed changes

coreneuron/gpu/nrn_acc_manager.cpp Show resolved Hide resolved

alkino approved these changes Dec 17, 2021

View reviewed changes

iomaganaris reviewed Dec 17, 2021

View reviewed changes

coreneuron/mechanism/mech/mod2c_core_thread.hpp Show resolved Hide resolved

coreneuron/permute/cellorder.cpp Outdated Show resolved Hide resolved

coreneuron/sim/scopmath/ssimplic_thread.cpp Show resolved Hide resolved

Apply suggestions from code review

9f8da7c

Co-authored-by: Ioannis Magkanaris <[email protected]>

olupton force-pushed the olupton/reenable-ci branch from 8954db4 to 9f8da7c Compare December 17, 2021 10:53

pramodk approved these changes Dec 17, 2021

View reviewed changes

olupton merged commit 3fc7037 into hackathon_main Dec 17, 2021

olupton deleted the olupton/reenable-ci branch December 17, 2021 13:53

Re-enable CI tests on hackathon branch #717

Re-enable CI tests on hackathon branch #717

Uh oh!

Conversation

olupton commented Dec 14, 2021 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

bbpbuildbot commented Dec 14, 2021

Uh oh!

bbpbuildbot commented Dec 14, 2021

Uh oh!

bbpbuildbot commented Dec 14, 2021

Uh oh!

olupton commented Dec 14, 2021

Uh oh!

bbpbuildbot commented Dec 15, 2021

Uh oh!

bbpbuildbot commented Dec 15, 2021

Uh oh!

bbpbuildbot commented Dec 16, 2021

Uh oh!

Uh oh!

iomaganaris left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

olupton commented Dec 17, 2021 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

iomaganaris commented Dec 17, 2021

Uh oh!

pramodk left a comment

Choose a reason for hiding this comment

Uh oh!

pramodk Dec 17, 2021

Choose a reason for hiding this comment

Uh oh!

olupton Dec 17, 2021

Choose a reason for hiding this comment

Uh oh!

bbpbuildbot commented Dec 17, 2021

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

5 participants

olupton commented Dec 14, 2021 •

edited

Loading

olupton commented Dec 17, 2021 •

edited

Loading