Software¶
Known Issues
Please see the Known Issues section at the bottom of this page regarding known software problems and incompatibilities on the Cori GPU nodes. If you encounter an issue which is not documented here, please file a ticket at the NERSC Help Desk, selecting 'Cori GPU' as the 'Resource' in the ticket.
Note about cross-compiling
Nearly all software provided by Cray (cray-petsc, cray-fftw, cray-hdf5,
etc.) is not usable on the Cori GPU nodes. This is because the GPU nodes have
different hardware and run a different OS. Only a select subset of modules
available on Cori are designed to work on the GPU nodes.
This means you will likely need to compile your own software directly on the GPU nodes themselves, rather than cross-compiling for the GPU nodes on a login node.
The best way to access a GPU node using modules designed to work on the GPU
nodes is to purge your default modules first, then load esslurm and the other
GPU modules you need, and then request the nodes, e.g.,
user@cori02:~> module purge && module load esslurm cuda pgi mvapich2 user@cori02:~> salloc -C gpu -t 60 -N 1 -c 10 --gres=gpu:1 -A <account> salloc: Granted job allocation 12345 salloc: Waiting for resource configuration salloc: Nodes cgpu12 are ready for job user@cgpu12:~>
MPI¶
Both MVAPICH2 and OpenMPI with UCX support are provided on the GPU nodes. Details about each are provided below.
MVAPICH2¶
MVAPICH2 is available via the mvapich2 module. It supports three compilers:
- GCC (via the
gccmodule) - PGI (via the
pgimodule) - Intel (via the
intelmodule)
The mvapich2 module must be loaded after a compiler module and a cuda
module. Thus, to load MVAPICH2 with PGI:
module load pgi module load cuda module load mvapich2
or, more succinctly,
module load pgi cuda mvapich2
To load MVAPICH2 with GCC or Intel support, replace pgi in this example with
gcc or intel.
After the mvapich2 module is loaded, the MPI compiler wrappers will be
available as mpicc, mpic++, and mpif90.
Cross-compiling with mvapich2 module from Cori login nodes does not work
Attempting to cross-compile a code on the Cori login nodes using the
mvapich2 compiler wrappers will result in an error like the following:
/global/common/cori/software/mvapich2/2.3/pgi/18.10/lib/libmpi.so: undefined reference to `ibv_reg_xrc_rcv_qp@IBVERBS_1.1' /global/common/cori/software/mvapich2/2.3/pgi/18.10/lib/libmpi.so: undefined reference to `ibv_close_xrc_domain@IBVERBS_1.1' /global/common/cori/software/mvapich2/2.3/pgi/18.10/lib/libmpi.so: undefined reference to `ibv_unreg_xrc_rcv_qp@IBVERBS_1.1' /global/common/cori/software/mvapich2/2.3/pgi/18.10/lib/libmpi.so: undefined reference to `ibv_open_xrc_domain@IBVERBS_1.1' /global/common/cori/software/mvapich2/2.3/pgi/18.10/lib/libmpi.so: undefined reference to `ibv_modify_xrc_rcv_qp@IBVERBS_1.1' /global/common/cori/software/mvapich2/2.3/pgi/18.10/lib/libmpi.so: undefined reference to `ibv_create_xrc_rcv_qp@IBVERBS_1.1' /global/common/cori/software/mvapich2/2.3/pgi/18.10/lib/libmpi.so: undefined reference to `ibv_create_xrc_srq@IBVERBS_1.1'
To avoid this error, one must invoke the mvapich2 compiler wrappers
directly on a Cori GPU node.
MPI_THREAD_MULTIPLE with MVAPICH2
By default, a code compiled with MVAPICH2 will support only
MPI_THREAD_SINGLE, even if a higher threading model is requested in
MPI_Init_thread(). This is by design; see
this
page for more information. If one requests a higher level of threading support, one will encounter the following runtime warning:
user@cgpu04:~/> mpicc -o main.ex main.c user@cgpu04:~/> srun -n 2 -c 2 ./main.ex Requested MPI_THREAD_MULTIPLE, got MPI_THREAD_SINGLE Hello world from processor cgpu04, rank 0 out of 2 processors Hello world from processor cgpu04, rank 1 out of 2 processors Error in system call pthread_mutex_destroy: Device or resource busy src/mpi/init/initthread.c:241 Error in system call pthread_mutex_destroy: Device or resource busy src/mpi/init/initthread.c:241
MPI_THREAD_MULTIPLE,
one must disable MVAPICH2's default task binding behavior by setting the
environment variable MV2_ENABLE_AFFINITY=0 during execution:
user@cgpu04:~> MV2_ENABLE_AFFINITY=0 srun -n 2 -c 2 ./main.ex Requested MPI_THREAD_MULTIPLE, got MPI_THREAD_MULTIPLE Hello world from processor cgpu04, rank 0 out of 2 processors Hello world from processor cgpu04, rank 1 out of 2 processors
OpenMPI + UCX¶
OpenMPI with UCX support is also installed as the openmpi/4.0.1-ucx-1.6
module. It supports only GCC 7.3.0 and PGI 19.7.
GPU Software Support¶
There are many different ways to offload code to GPUs. We provide software support for several of these methods on the GPU nodes.
CUDA¶
The CUDA SDK is available via the cuda modules. The SDK includes the nvcc
CUDA C/C++ compiler, the nvprof and NSight profiling tools, the cuda-gdb
debugger, and others.
Additionally, the LLVM/clang compiler is also a valid CUDA compiler. One can
replace the nvcc command from the CUDA SDK with clang
--cuda-gpu-arch=<arch>, where <arch> on the Cori GPU nodes is sm_70. If
using clang as a CUDA compiler, one usually will also need to add the
-I/path/to/cuda/include and -L/path/to/cuda/lib64 flags manually, since
nvcc includes them implicitly.
OpenMP¶
Several compilers have some support for OpenMP offloading to GPUs via the omp
target directive.
GCC¶
GCC 8.1.1 has some support for OpenMP offloading. This compiler is available
via the gcc/8.1.1-openacc-gcc-8-branch-20190215 module, which depends on the
cuda/9.2.148 module.
OpenMP offloading with gcc looks something like
gcc -fopenmp -foffload=nvptx-none="-Ofast -lm -misa=sm_35" base.c -c
LLVM/clang¶
The clang/clang++ LLVM compilers support GPU offloading with OpenMP. The 'raw' compilers are available via the following modules:
llvm/9.0.0-git_20190220llvm/10.0.0-git_20190828
or you can load the corresponding PrgEnv-llvm modules:
PrgEnv-llvm/9.0.0-git_20190220PrgEnv-llvm/10.0.0-git_20190828
which loads the appropriate LLVM, CUDA, and MVAPICH2 modules.
Enabling GPU offloading with OpenMP in the clang compiler looks like:
clang -fopenmp -fopenmp-targets=nvptx64-nvidia-cuda base.c -c
Using the clang++ compiler
The clang++ compiler will fail unless you add a compiler option to use an
official C++ standard, e.g. -std=c++11. The issue seems to be related to
GPU-offload support for GCC extensions, e.g. __float128 type.
Intrinsic math functions in GPU offloaded regions
The clang/clang++ compilers belonging to the llvm/9.0.0-git_20190220
module are unable to compile OpenMP target regions which call <math.h>
functions, e.g. log() and exp(). The compilers also incorrectly handle
OpenMP target regions inside static libraries -- your application will fail at
runtime when encountering the static library OpenMP target region. If you need
either of these capabilities please use the module
PrgEnv-llvm/10.0.0-git_20190828.
CCE¶
The Cray compilers ('CCE') have the most mature OpenMP offloading capabilities of any compiler on the Cori GPU nodes currently, especially amongst Fortran compilers. Cray does not officially supported CCE on the Cori GPU nodes, but it can be made to work by careful loading/unloading of modules.
# load the appropriate modules module load cdt/19.03 module swap PrgEnv-{intel,cray} module swap craype-{haswell,x86-skylake} module unload cray-libsci module load cudatoolkit craype-accel-nvidia70 # compile the code cc -h noacc -o my_openmp_code.ex my_openmp_code.c # C code CC -h noacc -o my_openmp_code.ex my_openmp_code.cpp # C++ code ftn -h noacc -o my_openmp_code.ex my_openmp_code.f90 # Fortran code
Do not module purge if using CCE
Unlike most other compilers and modules used on the Cori GPU nodes, which
should be preceded with module purge, the CCE compilers depend on the
default Cray module environment, and therefore one should not execute module
purge if one desires to use the CCE compilers.
You can add the flag -h list=a to any of the CCE compilers to produce an
optimization report (named <source_file>.lst) which indicates which regions
of the code were successfully offloaded to the GPU. For example, in the OpenMP
offload version of the
CloverLeaf benchmark code, the
Cray compiler outputs a diagnostic report for each source file which includes
sections such as:
%%% L o o p m a r k L e g e n d %%% Primary Loop Type Modifiers ------- ---- ---- --------- A - Pattern matched a - atomic memory operation b - blocked C - Collapsed c - conditional and/or computed D - Deleted E - Cloned F - Flat - No calls f - fused G - Accelerated g - partitioned I - Inlined i - interchanged M - Multithreaded m - partitioned n - non-blocking remote transfer p - partial R - Rerolling r - unrolled s - shortloop V - Vectorized w - unwound ... 105. #pragma omp target teams distribute if(halo_offload) 106. //#pragma omp parallel for 107. gG-----< for (int j = x_min-depth; j <= x_max+depth; j++) 108. gG { 109. gG #pragma ivdep 110. + gG r2--< for (int k = 1; k <= depth; k++) 111. gG r2 { 112. gG r2 density0[FTNREF2D(j ,y_max+k,x_max+4,x_min-2,y_min-2)] = 113. gG r2 density0[FTNREF2D(j ,y_max+1-k,x_max+4,x_min-2,y_min-2)]; 114. gG r2--> } 115. gG-----> } ... CC-6405 CC: ACCEL update_halo_kernel_c_, File = ext_update_halo.c, Line = 107 A region starting at line 107 and ending at line 115 was placed on the accelerator. CC-6430 CC: ACCEL update_halo_kernel_c_, File = ext_update_halo.c, Line = 107 A loop was partitioned across the threadblocks and the 128 threads within a threadblock.
OpenACC¶
Several compilers on the GPU nodes also support GPU offloading with OpenACC directives.
GCC¶
The GCC module available via gcc/8.1.1-openacc-gcc-8-branch-20190215 also
supports OpenACC offloading for GPUs. Invoking OpenACC looks like:
gcc -fopenacc -foffload=nvptx-none="-Ofast -lm -misa=sm_35" base.c -c
PGI¶
The PGI compilers support OpenACC offloading and are available via the pgi
modules.
Invoking OpenACC in the PGI compilers looks like:
pgc++ -acc -ta=tesla:cc70 base.c -c
Documentation for the PGI compiler is provided here.
CCE¶
The Cray compilers support OpenACC offloading, up to OpenACC v2.0. There is no planned support for newer versions of OpenACC.
module load cdt/18.12 module load PrgEnv-cray module unload cray-libsci module load cudatoolkit craype-accel-nvidia70 CC -h cpu=skylake -h noomp -h acc -h accel=nvidia70 base.c -c
Python¶
CuPy and Numba¶
Using Python on NVIDIA GPUs is commonly done via CuPy and/or Numba, although there are many other options including JAX, PyCUDA, PyOpenCL, and Legate.
In these docs we'll cover only Numba and CuPy. To use either or both of these
tools, you'll need to build your own conda environment. To do this, first make
sure that you have both the python and cuda modules loaded
module load python cuda
conda create -n my_gpu_env python=3.7 numpy numba
Don't conda install CuPy
Although you can conda install CuPy, you'll get a more up-to-date version if you do as the developers recommend and pip install it. For Cuda 10.2 (as of Feb 2020) you can install by:
pip install cupy-cuda102
conda uninstall cudatoolkit cudnn nccl
Note that your CuPy version must match your Cuda version. As of Feb 2020, the default version of Cuda on cori gpu is 10.2, hence you need CuPy 10.2.
Now that your my_gpu_env has both CuPy and Numba installed, you're ready to
start running Python on GPUs.
To summarize, here's what you need to do:
user@cgpu12:~> module load python cuda user@cgpu12:~> source activate my_gpu_env (my_gpu_env)user@cgpu12:~> srun python test_gpu.py
NVIDIA RAPIDS¶
NVIDIA RAPIDS is a tool that allows familiar Python libraries like scikit-learn
and Pandas to run on GPUs. We provide a RAPIDS kernel which you will find when
you log into jupyter.nersc.gov.
If you would like your own custom conda environment and/or Jupyter kernel that contains RAPIDS, you can follow the directions below.
1) Make sure you are on a Cori gpu node
2) Since RAPIDS only supports Cuda 10.1 as of March 2020:
module load python cuda/10.1.243
3) conda create -n rapids_env python=3.7
4) source activate rapids_env
5) Using the tool at https://rapids.ai/start.html, we generated the following command:
conda install -c rapidsai -c nvidia -c conda-forge \
-c defaults rapids=0.12 python=3.7 cudatoolkit=10.1
6) If you intend to use RAPIDS via scripts/command line, you're ready to go. If you
would like to create your own RAPIDS kernel to use in Jupyter, you'll need to
conda install ipykernel
and
python -m ipykernel install --user --name rapids_env --display-name rapids
7) You'll need to restart your Jupyter server. When you log in, you should now see
your rapids kernel as an option.
For more information about how to use NVIDIA RAPIDS, please see our Examples page.
MPI4py¶
You can build mpi4py and install it into a conda environment on Cori to be
used with one of the MPI implementations available for use with the GPU nodes.
First, request an interactive session on a GPU node:
user@cori02:~> module purge; module load python esslurm user@cori02:~> salloc -C gpu -A <account> -t 30 --gres=gpu:1 -c 10
Then, on the GPU node, create or activate a conda environment, load your MPI
implementation of choice (including relevant compiler), download mpi4py,
and build/install the software using the mpicc wrapper:
user@cgpu12:~> conda create -n mpi4pygpu python=2.7 user@cgpu12:~> source activate mpi4pygpu (mpi4pygpu) user@cgpu12:~> module load gcc/7.3.0 cuda mvapich2 (or pgi/intel instead of gcc) (mpi4pygpu) user@cgpu12:~> wget https://bitbucket.org/mpi4py/mpi4py/downloads/mpi4py-3.0.0.tar.gz (mpi4pygpu) user@cgpu12:~> tar zxvf mpi4py-3.0.0.tar.gz (mpi4pygpu) user@cgpu12:~> cd mpi4py-3.0.0 (mpi4pygpu) user@cgpu12:~> python setup.py build --mpicc=mpicc (mpi4pygpu) user@cgpu12:~> python setup.py install
Deep Learning Software¶
Tensorflow:
module load tensorflow/gpu-1.13.1-py36
PyTorch:
module load pytorch/v1.1.0-gpu
CUDA Fortran¶
The PGI Fortran compiler supports CUDA Fortran.
Compiler bugs
If you find bugs in the compilers (wrong answers, compiler crashing, etc.), PLEASE REPORT THEM TO NERSC! Any OpenMP target issues can be sent directly to Chris Daley: csdaley@lbl.gov. Many compilers are still in early phases of GPU enablement and depend on bug reports to fix these bugs quickly.
Intel MKL¶
To use routines provided by the Intel MKL, load one of the available intel
modules before compiling and running your code:
module load intel
Module names
Be sure to load an intel compiler module and not the Intel programming
environment module (PrgEnv-intel).
To determine the appropriate link lines for your code, use the Intel MKL Link Line Advisor.
Profiling¶
The Cori GPU nodes provide a few tools for profiling GPU code.
nvprof¶
nvprof has been CUDA's standard profiling tool for several years. It is easy to
use - one simply inserts the word nvprof in front of their application in the
srun command, and it will profile the code and generate a report:
user@cgpu17:~/samples/bin/x86_64/linux/release> srun -n 1 ./nvgraph_SpectralClustering GPU Device 0: "Tesla V100-SXM2-16GB" with compute capability 7.0 Modularity_score: 0.371466 Hit rate : 100.000000% (34 hits) Done! user@cgpu17:~/samples/bin/x86_64/linux/release> srun -n 1 nvprof ./nvgraph_SpectralClustering ==152717== NVPROF is profiling process 152717, command: ./nvgraph_SpectralClustering GPU Device 0: "Tesla V100-SXM2-16GB" with compute capability 7.0 Modularity_score: 0.371466 Hit rate : 100.000000% (34 hits) Done! ==152717== Profiling application: ./nvgraph_SpectralClustering ==152717== Profiling result: Type Time(%) Time Calls Avg Min Max Name GPU activities: 26.12% 22.309ms 6162 3.6200us 3.3270us 5.4400us void nrm2_kernel<float, float, float, int=1, int=0, int=128, int=0>(cublasNrm2Params<float, float>) 16.07% 13.722ms 8850 1.5500us 1.5040us 2.1120us [CUDA memcpy DtoH] 13.07% 11.161ms 8620 1.2940us 1.2160us 1.9520us void axpy_kernel_val<float, float, int=0>(cublasAxpyParamsVal<float, float, float>) 11.51% 9.8280ms 5752 1.7080us 1.6310us 13.600us void dot_kernel<float, float, float, int=128, int=0, int=0>(cublasDotParams<float, float>) 10.79% 9.2173ms 2877 3.2030us 3.1040us 13.344us void csrMv_kernel<float, float, float, int=128, int=1, int=2>(cusparseCsrMvParams<float, float, float>) 10.28% 8.7775ms 5752 1.5250us 1.4710us 2.4960us void reduce_1Block_kernel<float, float, float, int=128, int=7>(float*, int, float*) 5.54% 4.7290ms 3295 1.4350us 1.3760us 1.8240us [CUDA memcpy HtoD] 4.23% 3.6102ms 3079 1.1720us 1.1200us 1.4400us void scal_kernel_val<float, float, int=0>(cublasScalParamsVal<float, float>) 1.06% 908.79us 206 4.4110us 3.5190us 8.1920us volta_sgemm_128x32_nn 0.57% 487.16us 413 1.1790us 1.1520us 1.6640us [CUDA memcpy DtoD] ...
nvprof is part of the CUDA SDK (available via the cuda modules on the Cori
GPU nodes) but after CUDA 10 will be replaced by NSight Compute, NVIDIA's new
profiling tool.
Documentation about nvprof is here.
Using SLURM_TASK_ID with nvprof using multiple MPI ranks¶
nvprof profiles only one task at a time; if one profiles a GPU code which has
multiple tasks (e.g., multiple MPI ranks), nvprof will save one profile per
task if used with the -o flag. However, each file must have a unique
filename, or else all nvprof tasks will attempt to write to the same file,
typically resulting in an unusuable profiling result.
However, one can use the SLURM_TASK_PID environment variable to save each
nvprof profiling result to a unique file, which can then be combined to
generate a single result in the nvvp GUI. To do this, one must invoke a bash
shell inside the srun statement in order to return a unique value for the
SLURM_TASK_PID for each task:
user@cgpu02:~> srun -n 2 -c 2 bash -c 'echo $SLURM_TASK_PID' 36032 36031
bash -c statement, each task will return the same
value of SLURM_TASK_ID, because the wrong bash shell is interpreting the
value of the variable:
user@cgpu02:~> srun -n 2 -c 2 echo $SLURM_TASK_PID 35765 35765
The easiest way to use SLURM_TASK_PID to produce a unique profiling database
per MPI rank is to use it inside a shell script which is then invoked by
srun, e.g.,:
user@cgpu02:~> cat profile.sh #!/bin/bash file_prefix=$(bash -c 'echo $SLURM_TASK_PID') nvprof -o result_${file_prefix}.nvvp ./main.ex user@cgpu02:~> srun -n 2 -c 2 --cpu-bind=cores ./profile.sh ==44490== NVPROF is profiling process 44490, command: ./main.exe ==44492== NVPROF is profiling process 44492, command: ./main.exe ... (code runs) ... ==44490== Generated result file: /path/to/result/result_44468.nvvp ==44492== Generated result file: /path/to/result/result_44467.nvvp
.nvvp files in the nvvp GUI in order to
produce a unified profiling result of the entire application.
nvvp¶
nvvp is the profiling GPU which accompanies nvprof. It is used for displaying profiling information collected by nvprof in a GUI. Since X11 window forwarding via SSH is typically slow, one will enjoy a much better nvvp experience by running it on a Cori login node using NoMachine. Example output of nvvp is shown below:
NSight¶
NSight is NVIDIA's new profiling suite which will replace nvprof after CUDA 10. It measures much of the same information as nvprof, but organizes information in different ways. NSight is divided into two separate tools: NSight Compute, and NSight Systems.
NSight Systems¶
The way to achieve approximately the same output as nvprof ./my_code.exe is
to use nsys profile --stats=true ./my_code.exe, which is part of NSight
Systems. The output resembles the following:
Generating CUDA API Statistics... CUDA API Statistics (nanoseconds) Time(%) Total Time Calls Average Minimum Maximum Name ------- -------------- ---------- -------------- -------------- -------------- -------------------------------------------------------------------------------- 89.0 520526780 4002 130066.7 2119 1275014 cuStreamSynchronize 4.9 28856541 1 28856541.0 28856541 28856541 cuMemHostAlloc 2.8 16433738 3000 5477.9 4403 35531 cuLaunchKernel 0.9 5499511 1010 5445.1 1410 1306269 cuEventSynchronize 0.8 4816205 1005 4792.2 4170 22094 cuMemcpyDtoHAsync_v2 0.6 3420392 1000 3420.4 2884 20080 cuMemsetD32Async 0.4 2409796 5 481959.2 4670 1297188 cuMemAlloc_v2 0.3 1763291 1012 1742.4 1510 14416 cuEventRecord 0.2 1076739 1 1076739.0 1076739 1076739 cuMemAllocHost_v2 0.0 272120 1 272120.0 272120 272120 cuModuleLoadDataEx 0.0 75190 4 18797.5 15333 27138 cuMemcpyHtoDAsync_v2 0.0 12881 1 12881.0 12881 12881 cuStreamCreate 0.0 6902 4 1725.5 475 2795 cuEventCreate Generating CUDA Kernel Statistics... CUDA Kernel Statistics (nanoseconds) Time(%) Total Time Instances Average Minimum Maximum Name ------- -------------- ---------- -------------- -------------- -------------- -------------------------------------------------------------------------------- 56.3 284400677 1000 284400.7 274750 296989 laplace_85_gpu 43.2 218012722 1000 218012.7 215646 223102 laplace_102_gpu 0.5 2382122 1000 2382.1 2272 9280 laplace_85_gpu__red Generating CUDA Memory Operation Statistics... CUDA Memory Operation Statistics (nanoseconds) Time(%) Total Time Operations Average Minimum Maximum Name ------- -------------- ---------- -------------- -------------- -------------- -------------------------------------------------------------------------------- 49.2 6552324 1005 6519.7 1376 1295316 [CUDA memcpy DtoH] 40.9 5445551 4 1361387.7 1354995 1365236 [CUDA memcpy HtoD] 9.8 1308501 1000 1308.5 1279 1952 [CUDA memset] CUDA Memory Operation Statistics (KiB) Total Operations Average Minimum Maximum Name ------------------- -------------- ------------------- ----------------- ------------------- -------------------------------------------------------------------------------- 65539.906 1005 65.214 0.004 16383.937 [CUDA memcpy DtoH] 65536.000 4 16384.000 16384.000 16384.000 [CUDA memcpy HtoD] 3.906 1000 0.004 0.004 0.004 [CUDA memset]
NSight Systems will save the profiling output to a .qdrep file in the present
working directory. One can then view the .qdrep profiling database via the
NSight Systems GUI. Adding the --stats=true flag to nsys profile causes
NSight Systems to automatically convert the .qdrep file into a .sqlite
file. As with nvvp, the NSight Systems GUI is best viewed from a Cori login
node using NoMachine. Inside a NoMachine
session, one can launch the NSight Systems GUI with the nsight-sys. From the
GUI, one can import the .qdrep file to view the profiling output. An example
of this output is shown below:
NVIDIA documentation about NSight Systems is here.
NSight Compute¶
The NSight Compute tool enables deep dives into GPU code performance. The
command line interface to NSight Compute is nv-nsight-cu-cli. More
documentation about this tool is forthcoming.
NVIDIA documentation about NSight Compute is here.
How do I know if my code ran on the GPU?¶
While it is usually clear that a code has run at all, it is sometimes less clear whether the code ran on the CPU or the GPU. One way this ambiguity can arise is if one includes GPU offloading directives in the code, but does not use the appropriate compiler flag to enable those directives.
There are several ways to determine if your code actually ran on the GPU. It may be more useful to know when a code does not run on the GPU (especially when one expects that it should):
-
Run the code through an NVIDIA profiler such as NSight Compute or nvprof. If a code runs on the GPU, both profilers will print a summary following code execution:
An NSight Compute profile would look something like:==39359== Profiling application: ./laplace2d_acc ==39359== Profiling result: Type Time(%) Time Calls Avg Min Max Name GPU activities: 53.75% 264.48ms 1000 264.48us 254.33us 292.35us main_96_gpu 43.03% 211.73ms 1000 211.73us 210.27us 214.01us main_109_gpu 1.37% 6.7473ms 1004 6.7200us 1.2790us 1.3656ms [CUDA memcpy HtoD] 1.33% 6.5500ms 1005 6.5170us 1.4070us 1.2849ms [CUDA memcpy DtoH]
An NSight Compute profile of a code which does not run on the GPU at all will print the following message:
and an nvprof profile would print:user@cgpu12:~/tests> srun -n 1 nv-nsight-cu-cli ./a.out ==PROF== ERROR: Target application terminated before first instrumented API call. srun: error: cgpu12: task 0: Exited with exit code 255 srun: Terminating job step 123456.5 user@cgpu12:~/tests>
user@cgpu12:~/tests> srun -n 1 nvprof ./a.out ======== Warning: No profile data collected. user@cgpu12:~/tests>
-
Check compiler reports. If one writes a code with OpenACC directives and compiles it with the PGI compiler but does not include the flags needed to inform the compiler to use the directives, there will be no output:
Including the appropriate OpenACC flags to the compiler (in this caseuser@cori02:~> pgcc -I../common -Minfo=accel -o laplace2d_acc laplace2d.c user@cori02:~>
-acc -ta=nvidia) results in more output, including a note that the compiler generated Tesla code.user@cori02:~> pgcc -I../common -acc -ta=nvidia -Minfo=accel -o laplace2d_acc laplace2d.c main: 86, Generating copy(A[:][:]) Generating create(Anew[:][:]) 93, Loop is parallelizable 96, Loop is parallelizable Generating Tesla code 93, #pragma acc loop gang(32), vector(16) /* blockIdx.y threadIdx.y */ 96, #pragma acc loop gang(16), vector(32) /* blockIdx.x threadIdx.x */ 100, Generating implicit reduction(max:error) 106, Loop is parallelizable 109, Loop is parallelizable Generating Tesla code 106, #pragma acc loop gang, vector(4) /* blockIdx.y threadIdx.y */ 109, #pragma acc loop gang(16), vector(32) /* blockIdx.x threadIdx.x */ user@cori02:~>
Shifter with CUDA¶
Shifter works in a slightly different way on the Cori GPU nodes than it does on
the Haswell or KNL nodes. First, one should add the following to ENTRYPOINT
in the Shifter container:
export PATH=/opt/shifter/bin:${PATH} export LD_LIBRARY_PATH=/opt/shifter/lib:${LD_LIBRARY_PATH}
Next, one should load the cuda/shifter module; this will define the SHIFTER_CUDA_ROOT environment variable and point it to the version of the CUDA SDK installation which works in Shifter images.
cuda modules and Shifter images
Note that the normal cuda modules will not work inside Shifter containers.
Then one should invoke shifter from the job script as follows:
srun -n <num_task> -c <num_cpu> \ shifter \ --image=<your_image> \ --entrypoint \ --volume=${SHIFTER_CUDA_ROOT}:/opt/shifter:ro \ ./your_gpu_code.ex args
A complete is example is show below; this example is also provided in
$SHIFTER_CUDA_ROOT/example.
#!/bin/bash -e #SBATCH -A nstaff #SBATCH -N 1 #SBATCH -C gpu #SBATCH --gres=gpu:1 #SBATCH -t 00:10:00 #SBATCH --job-name=nvidia-shifter # provides SHIFTER_CUDA_ROOT module load cuda/shifter # # jrmadsen/tomopy:shifter container has the following in entrypoint: # # export PATH=/opt/shifter/bin:${PATH} # export LD_LIBRARY_PATH=/opt/shifter/lib:${LD_LIBRARY_PATH} # export TOMOPY_NUM_THREADS=32 export NUMEXPR_MAX_THREADS=80 export OMP_NUM_THREADS=1 srun -n 1 -c 1 \ shifter \ --image=jrmadsen/tomopy:shifter \ --entrypoint \ --volume=${SHIFTER_CUDA_ROOT}:/opt/shifter:ro \ python ${SHIFTER_CUDA_ROOT}/example/phantom.py -a mlem -i 100 -n 1 -p shepp2d -f png -s 256
Known Issues¶
MPS disabled indefinitely¶
NVIDIA's Multi-Process Service (MPS) enables multiple processes (typically MPI ranks) to execute kernels on a single GPU simultaneously. MPS can enable an application to achieve higher performance when a single process is unable to saturate the GPU's resources.
Unfortunately, a security vulnerability in the V100 GPUs was disclosed in February 2019 (see here, here, and here for more information) which renders data in GPU memory exposed to a side-channel attack if the GPU is accessed by multiple processes simultaneously. The vulnerability is not present if a GPU is allocated exclusively to a single process.
As a result of this security risk, NERSC has disabled MPS on Cori GPU until a mitigation for this vulnerability is implemented.
PGI 19.5 requires CUDA <= 10.1.105¶
The PGI v19.5 compiler using OpenACC directives is compatible with CUDA modules
only up to cuda/10.1.105. It is not compatible with cuda/10.1.168. If one
attempts to compile OpenACC code with the cuda/10.1.168 module loaded, one
encounters an error at runtime:
user@cgpu01:~/tests/OpenACC/vector_add> module list -l - Package -----------------------------+- Versions -+- Last mod. ------ Currently Loaded Modulefiles: esslurm 2019/02/08 22:01:04 pgi/19.5 2019/07/19 22:04:32 modules/3.2.10.6 2017/04/27 21:50:33 cuda/10.1.168 2019/07/19 22:01:27 user@cgpu01:~/tests/OpenACC/vector_add> pgf90 -acc -ta=tesla -o vector_add.ex vector_add.f90 user@cgpu01:~/tests/OpenACC/vector_add> srun -n 1 ./a.out Failing in Thread:0 call to cuInit returned error -1: Other srun: error: cgpu01: task 0: Exited with exit code 1 srun: Terminating job step 182246.6 user@cgpu01:~/tests/OpenACC/vector_add>
MVAPICH2 ptmalloc warnings with Python MPI codes¶
When running some MPI-enabled Python codes compiled with MVAPICH2, one may encounter the following warning:
WARNING: Error in initializing MVAPICH2 ptmalloc library.Continuing without InfiniBand registration cache support.
This is due to a bad interaction between MVAPICH's ptmalloc library and the
memory allocator used in Python. Details about this warning are provided
here.
As described in that page, one workaround is to set the LD_PRELOAD
environment variable:
export LD_PRELOAD=$MVAPICH2_DIR/lib/libmpi.so