[Hackathon] Basic OpenACC -> OpenMP migration. #693

olupton · 2021-11-23T08:59:38Z

Description
Migrate some offloaded kernels from OpenACC to OpenMP.

The idea with the macros defined here is that the macros nrn_acc_pragma and nrn_omp_pragma are used to annotate expressions that can be used with either OpenACC or OpenMP. i.e.

nrn_acc_pragma(atomic update)
nrn_omp_pragma(atomic update)
rhs += p * rhs;

will expand to either

#pragma acc atomic update
rhs += p * rhs;

or

#pragma omp atomic update
rhs += p * rhs;

Some other directives may be needed, for example if we want OpenACC directives in a build that prefers OpenMP offload then we may need code like

#ifdef CORENRN_PREFER_OPENMP_OFFLOAD
// Make sure OpenACC work is done before we start OpenMP work
if (nt->compute_gpu) {
    _Pragma("acc wait(nt->stream_id)")
}
#endif

Use certain branches for the GitLab/SimulationStack CI
CI_BRANCHES:NMODL_BRANCH=hackathon_main,NEURON_BRANCH=master,

bbpbuildbot · 2021-11-23T10:51:15Z

Logfiles from GitLab pipeline #26098 (:white_check_mark:) have been uploaded here!

Status and direct links:

bbpbuildbot · 2021-11-29T10:31:16Z

Logfiles from GitLab pipeline #27003 (:white_check_mark:) have been uploaded here!

Status and direct links:

bbpbuildbot · 2021-11-29T15:00:26Z

Logfiles from GitLab pipeline #27095 (:white_check_mark:) have been uploaded here!

Status and direct links:

bbpbuildbot · 2021-11-29T16:37:51Z

Logfiles from GitLab pipeline #27122 (:no_entry:) have been uploaded here!

Status and direct links:

bbpbuildbot · 2021-11-30T13:52:37Z

Logfiles from GitLab pipeline #27287 (:no_entry:) have been uploaded here!

Status and direct links:

bbpbuildbot · 2021-11-30T16:06:52Z

Logfiles from GitLab pipeline #27326 (:no_entry:) have been uploaded here!

Status and direct links:

bbpbuildbot · 2021-12-01T08:31:21Z

Logfiles from GitLab pipeline #27373 (:white_check_mark:) have been uploaded here!

Status and direct links:

bbpbuildbot · 2021-12-01T11:02:41Z

Logfiles from GitLab pipeline #27419 (:no_entry:) have been uploaded here!

Status and direct links:

bbpbuildbot · 2021-12-01T12:55:49Z

Logfiles from GitLab pipeline #27444 (:no_entry:) have been uploaded here!

Status and direct links:

coreneuron/gpu/nrn_acc_manager.cpp

So far: - Pass -mp=gpu when we pass -acc - Pass -gpu=lineinfo for better debug information. - Pass -Minfo=accel,mp for better compile time diagnostics. - Add nrn_{acc,omp}_pragma macros to make clang-format less painful. - Add omp_set_default_device call so the CTest suite works. - Transform one loop in the matrix solver from OpenACC to OpenMP. - Drop cc60 because of OpenMP offload incompatibility.

bbpbuildbot · 2021-12-02T11:03:17Z

Logfiles from GitLab pipeline #27587 (:no_entry:) have been uploaded here!

Status and direct links:

pramodk

Overall looks good to me. As discussed offline, we will fix the clang-formatted part of pragmas manually.

CMake/OpenAccHelper.cmake

coreneuron/gpu/nrn_acc_manager.cpp

coreneuron/io/lfp.cpp

coreneuron/mechanism/eion.cpp

coreneuron/mechanism/register_mech.cpp

coreneuron/permute/cellorder.cpp

coreneuron/sim/fadvance_core.cpp

These OpenACC pragmas are debug-only and have not been re-implemented for OpenMP target offload.

bbpbuildbot · 2021-12-06T12:27:24Z

Logfiles from GitLab pipeline #28163 (:no_entry:) have been uploaded here!

Status and direct links:

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

Logfiles from GitLab pipeline #28256 (:no_entry:) have been uploaded here!

Status and direct links:

bbpbuildbot · 2021-12-06T17:42:49Z

Logfiles from GitLab pipeline #28269 (:no_entry:) have been uploaded here!

Status and direct links:

bbpbuildbot · 2021-12-06T20:21:22Z

Logfiles from GitLab pipeline #28280 (:white_check_mark:) have been uploaded here!

Status and direct links:

Prefer CORENEURON_ prefixes for macros, CORENRN_ prefixes for CMake variables.

bbpbuildbot · 2021-12-07T08:31:25Z

Logfiles from GitLab pipeline #28331 (:no_entry:) have been uploaded here!

Status and direct links:

bbpbuildbot · 2021-12-07T09:28:52Z

Logfiles from GitLab pipeline #28336 (:white_check_mark:) have been uploaded here!

Status and direct links:

pramodk

LGTM. Only my open question is about seq part in cellorder.cpp. We can merge this and review it later if you like.

coreneuron/permute/cellorder.cpp

coreneuron/network/partrans.cpp

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/basic-openmp branch from e3c8352 to 713b7ad Compare November 29, 2021 09:25

olupton force-pushed the olupton/basic-openmp branch from 174211a to 80da782 Compare November 29, 2021 14:13

olupton closed this Nov 30, 2021

olupton reopened this Nov 30, 2021

olupton force-pushed the olupton/basic-openmp branch from 1effa13 to e868af4 Compare November 30, 2021 15:39

olupton closed this Dec 1, 2021

olupton reopened this Dec 1, 2021

alkino reviewed Dec 1, 2021

View reviewed changes

coreneuron/gpu/nrn_acc_manager.cpp Show resolved Hide resolved

olupton added 13 commits December 2, 2021 11:06

Simplify unified memory logic.

57f12be

nrn_{acc,omp}_pragma -> nrn_pragma_{acc,omp}.

9f14e63

Add --gpu to test.

4d67bf0

Default (BB5-valid) CORENRN_EXTERNAL_BENCHMARK_DATA.

5d9c7e7

Remove cuda_add_library.

385a34b

Define nrn_pragma_{acc,omp} in header.

8c6210e

Update NMODL with codegen fixes.

424b14f

Migrate more pragmas.

4c86be3

Move input data.

4ffb0bc

Migrate more.

388fe26

Migrate more directives.

e36f6ad

Remove more OpenACC from the main simulation section.

8cad9eb

olupton force-pushed the olupton/basic-openmp branch from 572066f to 8cad9eb Compare December 2, 2021 10:06

pramodk reviewed Dec 3, 2021

View reviewed changes

olupton added 6 commits December 6, 2021 11:16

Use nrn_pragma_acc for cell_permute=0 on GPU.

5c2a98c

These OpenACC pragmas are debug-only and have not been re-implemented for OpenMP target offload.

OpenMP target offload for fast_imem.

fb1139e

Don't print number of GPUs when quiet.

ee3588d

OpenMP target offload for partrans.

5580951

OpenMP target offload for finitialize.

a62b2ad

Fixup for nrnthread_v_transfer + OpenMP.

6d06c02

fix one weird clang-formatting in the nrn_acc_manager.cpp

e548dbf

olupton added 2 commits December 6, 2021 17:36

Set OMP_NUM_THREADS=1 for lfp_test.

184cdd7

Address review comments.

587ff37

Drop nowait and depend clauses.

c951423

olupton added 2 commits December 7, 2021 09:13

Rename CORENRN_PREFER_OPENMP_OFFLOAD.

6cb1b26

Prefer CORENEURON_ prefixes for macros, CORENRN_ prefixes for CMake variables.

Update NMODL with OpenMP codegen.

aefa0d4

olupton closed this Dec 7, 2021

olupton reopened this Dec 7, 2021

pramodk approved these changes Dec 7, 2021

View reviewed changes

coreneuron/permute/cellorder.cpp Show resolved Hide resolved

coreneuron/network/partrans.cpp Show resolved Hide resolved

olupton merged commit 21dc2c8 into hackathon_main Dec 7, 2021

olupton deleted the olupton/basic-openmp branch December 7, 2021 12:13

[Hackathon] Basic OpenACC -> OpenMP migration. #693

[Hackathon] Basic OpenACC -> OpenMP migration. #693

Uh oh!

Conversation

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

Uh oh!

Uh oh!

bbpbuildbot commented Nov 23, 2021

Uh oh!

bbpbuildbot commented Nov 29, 2021

Uh oh!

bbpbuildbot commented Nov 29, 2021

Uh oh!

bbpbuildbot commented Nov 29, 2021

Uh oh!

bbpbuildbot commented Nov 30, 2021

Uh oh!

bbpbuildbot commented Nov 30, 2021

Uh oh!

bbpbuildbot commented Dec 1, 2021

Uh oh!

bbpbuildbot commented Dec 1, 2021

Uh oh!

bbpbuildbot commented Dec 1, 2021

Uh oh!

Uh oh!

bbpbuildbot commented Dec 2, 2021

Uh oh!

pramodk left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

bbpbuildbot commented Dec 6, 2021

Uh oh!

bbpbuildbot commented Dec 6, 2021

Uh oh!

bbpbuildbot commented Dec 6, 2021

Uh oh!

bbpbuildbot commented Dec 6, 2021

Uh oh!

bbpbuildbot commented Dec 7, 2021

Uh oh!

bbpbuildbot commented Dec 7, 2021

Uh oh!

pramodk left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

4 participants

olupton commented Nov 23, 2021 •

edited

Loading