17.7 Updates and Additions

• PGI Accelerator Compilers
  • Added support for Tesla V100 GPUs in PGI OpenACC and CUDA Fortran. Based on the new NVIDIA Volta GV100 GPU, Tesla V100 offers more memory bandwidth, more streaming multiprocessors, next generation NVLink and new microarchitectural features that add up to better performance and programmability. For OpenACC and CUDA Fortran programmers, Tesla V100 offers improved hardware support and performance for CUDA Unified Memory features on both x86-64 and OpenPOWER processor-based systems. Use the sub-option cc70 with the -⁠ta=tesla or -⁠Mcuda compile- and link-time options to generate code for the Tesla V100 GPUs; the CUDA 9 toolkit must be used when targeting Volta GPUs.
  • Added initial support for the CUDA 9 toolkit. Requirements for using CUDA 9 with the PGI 17.7 compilers and profiler:
    • Install the release candidate (RC) or early access (EA) version of the CUDA 9 toolkit.
    • Specify the location of the CUDA 9 toolkit using the compiler option CUDAROOT; for example, if CUDA 9 is installed in /opt/cuda-9.0, add CUDAROOT=/opt/cuda-9.0 when compiling, linking, or invoking the profiler.
    • Use the sub-option cuda9.0 with the -⁠ta=tesla or -⁠Mcuda compile- and link-time options.
  • Added support for CUDA Unified Memory. PGI 17.7 compilers leverage Pascal and Volta GPU hardware features, NVLink, and CUDA Unified Memory to simplify OpenACC and CUDA Fortran programming on GPU-accelerated x86-64 and OpenPOWER processor-based servers. When CUDA Fortran or OpenACC allocatable data is placed in CUDA Unified Memory, no explicit data movement or data directives are needed. This simplifies GPU acceleration of applications that make extensive use of allocatable data, and allows you to focus on parallelization and scalability of your algorithms. Use the -⁠ta=tesla:managed compiler option to enable this feature. In previous PGI releases, compiler support for CUDA Unified Memory was provided as an optionally-installed, evaluation feature.
  • Added a pool allocator to minimize the number of calls to cudaMallocManaged() for CUDA Unified Memory support. The pool allocator is enabled by default when -⁠ta=tesla:managed or -⁠ta=tesla:pinned is used.
  • Added Beta support for automatic deep copy of Fortran derived types. This feature allows you to port applications with modern deeply nested data structures to Tesla GPUs using OpenACC. The PGI 17.7 compilers allow you to list aggregate Fortran data objects in OpenACC copy, copyin, copyout and update directives to move them between host and device memory including traversal and management of pointer-based objects within the aggregate data object. When enabled, full deep copy ensures that, when moving Fortran variables of derived type from host to device or device to host, the entire data structure, including pointers and allocatable arrays, is copied back and forth between host and device, or device and host, memory. To enable deep copy, use the deepcopy sub-option to -⁠ta=tesla. Two things to note: polymorphic data types are not supported, and the presence of overlapping pointers may cause runtime errors.
  • The cuSOLVER library is based on the cuBLAS and cuSPARSE libraries, and supports matrix factorization and triangular solve routines for dense matrices, a sparse least-squares solver and an eigenvalue solver, and provides a refactorization library useful for solving sequences of matrices with a shared sparsity pattern. You can now call optimized cuSolverDN routines from CUDA Fortran and OpenACC Fortran using the PGI-supplied interface module and the PGI-compiled version of the cuSOLVER libraries bundled with PGI 17.7. This same cuSolver library is callable from PGI OpenACC C/C⁠+⁠+ as they are built using PGI compilers and are compatible with and use the PGI OpenMP runtime.
• All PGI Compilers
  • Added a Beta LLVM code generator and OpenMP runtime to the PGI Linux/x86-64 compilers. The PGI Beta compilers are available as a separate download and installation package. While there are some limitations relative to the production PGI compilers, the Beta release includes support for all OpenACC, CUDA Fortran and OpenMP 4.5 features and delivers significant performance improvements on many C⁠+⁠+ applications. For more details, see Beta LLVM Code Generator.
  • Added initial support for OpenMP 4.5 syntax and features. The PGI Fortran, C, and C⁠+⁠+ compilers will now compile most OpenMP 4.5 programs for parallel execution across all the cores of a multicore CPU or server. target regions are implemented with default support for the multicore host as the target, and parallel and distribute loops are parallelized across all OpenMP threads. This feature is supported on Linux/x86 platforms with the Beta LLVM code generator only.
  • Improved inlining with the -⁠Minline option. You may see different functions inlined with this release than you observed in prior releases. In some cases, compilation may take noticeably longer to complete because of an increase in the number of functions inlined. Added the smallsize:number sub-option tomanage -⁠Minline which you can use to always inline functions of size smaller than number regardless of other size limits.
  • Added a new feature for program exploration called Program Analysis Summary Output (PASO). PASO enables the exporting of information gathered from the internal representation (IR) of various compilation stages in an organized, human readable format. This data forms a detailed summary of the program's internal structure, spanning not just one but all the files that make up a project. To enable PASO, compile and link with the -⁠Msummary option. To find out more about how PASO works and how you might leverage it, refer to Program Analysis Using Summary Output. PASO is available on all platforms supported by PGI except macOS.
• C⁠+⁠+ Compiler
  • Improved the raw performance on the LCALS loops benchmarks by an average of 20%, as well as eliminated the lambda abstraction penalty associated with these loops.
  • Added support for lambdas with capture in OpenACC GPU-accelerated compute regions.
  • Added incremental support for C⁠+⁠+⁠17 features consistent with EDG 4.13. This support includes:
    • new rules for auto deduction from braced-init-list
    • allow typename in template template parameter list
    • attributes for namespaces and enumerators
    • u8 character literals
    • nested namespace definitions
    • extended static_assert
    • generalized range-based for
    • remove deprecated register keyword
    • remove deprecated operator++(bool)
    • __has_include
    • [[fallthrough]]
    • [[nodiscard]]
    • [[maybe_unused]]
    • hexadecimal floating literals
    • using attribute namespaces without repetition
    • standard and non-standard attributes
  • Added support for interoperability with GNU 6.3.
  • Added support for the PGI C++ compilers as a CUDA 9.0 NVCC host compiler.
• Profiler

  The following new profiler features are enabled only when invoking the profiler with CUDA 9. To use CUDA 9 with PGI 17.7, start the profiler with the CUDAROOT option set to the CUDA 9 installation location. For example, if CUDA 9 is installed in /opt/cuda-9.0, add CUDAROOT=/opt/cuda-9.0 when launching the profiler.

  • Enhanced unified memory profiling:
    • Added association between unified memory events/CPU page faults and the source location of memory allocation.
    • Added new unified memory profiling events for page thrashing, throttling and remote map.
    • Added support for switching between segment and non-segment timelines.
    • Added filtering of events based on the virtual address, migration reason or page fault access type.
  • Added an OpenACC details table summarizing each recorded event.
  • Added support for the profiling of cooperative kernel launches.
  • Added NVLink events to the timeline.
  • Added memory throughput to the NVLink topology diagram.
  • Increased the options available for multi-hop remote profiling.
  • Added support for OpenACC profiling on all multicore systems.
• Documentation

  Added HTML-formatted documentation for all PGI compilers and tools online at https://www.pgroup.com. The new format should make it easier to use any Internet-connected system or device to search or reference PGI documentation. PDF versions of all documents also remain available.

• Operating Systems
  • Added support for Fedora 25 and openSUSE 13.2.

17.5 Updates and Additions

• CUDA Toolkit
  • This version includes an updated nvlink linker utility that addresses an issue with the link ordering of object files on Linux.
• Operating Systems
  • Added support for Ubuntu 17.04 and RHEL 6.9 Linux distributions.

17.4 Updates and Additions

• PGI Accelerator Compilers
  • Added support for atomic add and atomic subtract in device code for Fortran's single precision complex data types.
  • Added device versions of isfinite, isnan, isinf, and round.
  • Added support for CUDA Fortran warp-vote operations.
  • Changed the impact of -⁠g on the optimization level used when compiling device code. For host code, -⁠g sets the compiler's optimization level to zero unless another optimization level is specified using a -⁠O option. For device code, the same is now true. For details about generating debug information for device code, refer to the PGI Compiler Reference Guide's section on DWARF Debugging Formats.
  • Updated the version of nvlink to 8.0.73; this version of the device linker allows an arbitrary ordering of object files on its command line.
• PGI Tools

  Improved debugging support of C⁠+⁠+ programs on macOS versions 10.11 (El Capitan) and 10.12 (Sierra).

17.3 Updates and Additions

17.1 Updates and Additions

• All PGI Compilers
  • Updated how floating point divides are computed to guarantee results will be uniform for scalar and vector operations. This change can cause numerical differences between PGI 17.1 and previous PGI releases. In rare cases, this change can cause an increase in execution time.
  • Improved inlining with the -⁠Minline option. You may see different functions inlined with this release than you observed in prior releases. In some cases, compilation may take longer to complete because of an increase in the number of functions inlined. Some of the -⁠Minline sub-options, which you can use to control inlining, have also changed from previous releases:
    • Added totalsize:n to limit inlining to the total size of n where n is the size of the combined program units on a per file basis.
    • Changed size:n to maxsize:n which allows inlining only of functions smaller than approximately n lines. The compilers silently convert the previous size:n to maxsize:n.
    • Dropped levels:n which limited inlining to n levels of functions. The compilers silently ignore levels:n.
  • Added the -⁠cpp option as an alias for the -⁠Mpreprocess option.
• PGI C⁠+⁠+ Compiler
  • Integrated EDG release 4.13 and enabled interoperability with GCC 5.2, 5.3, 5.4, 6.0, 6.1, 6.2, and 6.3. To take full advantage of the new C⁠+⁠+⁠11 features supported by GCC 5.1 and later, use the PGC⁠+⁠+ compiler option --c⁠+⁠+⁠11.
  • Comprehensive support for C⁠+⁠+⁠14; co-installation with GCC version 5.0 or greater is required. Enable C⁠+⁠+⁠14 features with the PGC⁠+⁠+ compiler option --c⁠+⁠+⁠14.
  • Added C⁠+⁠+⁠11 support to PGC⁠+⁠+⁠ when PGC⁠+⁠+ is used as the nvcc host compiler; this support requires a patched post-CUDA-8.0-production version of nvcc. Contact sales@pgroup.com for more information.
  • Improved exception handling support and dropped the following related compiler options: -⁠-⁠[no_]exceptions, -⁠-⁠need_exception_spec, -⁠-⁠sjlj_eh, -⁠-⁠[no]zc_eh.
  • Added -⁠M[no]variadic_macros to allow or disallow variadic macros. Variadic macros are allowed by default.
  • Deprecated the compiler options -⁠-⁠edg and -⁠-⁠gnu.
  • GNU changed the GCC STL header files in version 5.0. These changes may introduce compatibility issues with code compiled with earlier GCC versions. The GCC-supplied workaround to support compatibility with pre-GCC 5.0 compiled objects and libraries is to compile with -⁠D_GLIBCXX_USE_CXX11_ABI=0. For more information, refer to https://gcc.gnu.org/onlinedocs/libstdc++/manual/using_dual_abi.html.
• PGI Accelerator OpenACC Compilers
  • Dropped support for CUDA 7.0.
  • Changed the default version of the CUDA Toolkit used by the compilers from CUDA 7.0 to 7.5. The CUDA Toolkit 8.0 can be used instead of 7.5 by adding the sub-option cuda8.0 to the -⁠ta=tesla or -⁠Mcuda compile- and link-time options.
  • Removed compute capability 2.0 (Fermi) from the compilers' default set of compute capabilities. The cc20 sub-option to the -⁠ta=tesla and -⁠Mcuda options is deprecated.
  • Improved data management within the OpenACC cache directive.
  • Added support for additional OpenACC 2.5 features:
    • Changed the behavior of the exit data directive to decrement the dynamic reference count.
    • Added the new optional finalize clause to set the dynamic reference count to zero.
    • Added the if_present clause to the update directive which changes the behavior when data is not present from a runtime error to a no-op.
    • Added new init, shutdown, and set directives.
    • Added new API routines to get and set the default async queue value.
    • Added support for the new definition of routine bind clause.
    • Updated the value of _OPENACC to 201510.
    With the exception of nested parallelism, declare link, and adding restrictions to cache clause variable references to variables within a cached region, OpenACC 2.5 feature support is complete in PGI 2017.
• PGI Tools
  • Updated the version of the JRE included in all installation packages to version 1.8.0_112.
  • Added variable rollover to the PGI debugger's graphical interface.
  • The PGI debugger requires libncurses.so.5. Some newer Linux distributions, such as Fedora 24, do not include this version of this library by default. In this situation, install the ncurses compatibility library to use the debugger.
• Libraries
  • PGI products for Linux ship with Open MPI 1.10.2.
  • Users with PGI subscription support can access separate downloads of MVAPICH 2.2 and MPICH 3.2.
  • PGI-built versions of NetCDF 4.4.1.1 and Parallel NetCDF 1.7.0 are available online at http://www.pgroup.com/resources/tips.htm
  • PGI-built versions of the Earth System Modeling Framework (ESMF) 7.0.0, one per PGI-built MPI distribution, are available online at the address above.
• Other Features and Additions
  • Added support for these new Linux x86-64 operating systems: Fedora 24, RHEL 7.3, Ubuntu 16.10.
  • Added compiler and profiler support for macOS version 10.12 (Sierra); the PGI debugger is not supported on Sierra in PGI 17.1.
  • Added support for Microsoft Windows Server 2016.
• Deprecations and Eliminations
  • PGI 2017 supports 64-bit operating systems only. Compiling 32-bit applications for execution on either 32-bit or 64-bit operating systems is no longer supported on any platform.
  • The PGI 2017 release for macOS no longer supports CUDA Fortran and OpenACC running on GPUs or CUDA-x86 running on CPUs. OpenACC targeting multi-core CPUs is still supported using the -⁠ta sub-options host and multicore. The -⁠acc option now implies -⁠ta=multicore on macOS.
  • Dropped support for the CUDA 7.0 toolkit.
  • Dropped support for OS X 10.7 (Lion) and 10.8 (Mountain Lion).

CUDA Unified Memory

In the PGI 17.7 release, the use of CUDA Unified Memory for allocatable data moved from a Beta feature to production. This feature, described in detail in the OpenACC and CUDA Unified MemoryPGInsider article, is available with the Linux/x86-64 and Linux/OpenPOWER compilers. It is supported on Linux/x86-64 using both the default PGI code generator and the Beta LLVM-based code generator. To enable this feature, add the option -⁠ta=tesla:managed to the compiler and linker command lines.

In the presence of -⁠ta=tesla:managed, all C/C⁠+⁠+/Fortran explicit allocation statements in a program unit are replaced by equivalent "managed" data allocation calls that place the data in CUDA Unified Memory. Managed data share a single address for CPU/GPU and data movement between CPU and GPU memories is implicitly handled by the CUDA driver. Therefore, OpenACC data clauses and directives are not needed for "managed" data. They are essentially ignored, and in fact can be omitted.

When a program allocates managed memory, it allocates host pinned memory as well as device memory thus making allocate and free operations somewhat more expensive and data transfers somewhat faster. A memory pool allocator is used to mitigate the overhead of the allocate and free operations. The pool allocator is enabled by default for -⁠ta=tesla:managed or -⁠ta=tesla:pinned. In the PGI 17.7 release, the presence of -⁠Mcuda disables the pool allocator; we are working on lifting that restriction in an upcoming release.

Data movement of managed data is controlled by the NVIDIA CUDA GPU driver; whenever data is accessed on the CPU or the GPU, it could trigger a data transfer if the last time it was accessed was not on the same device. In some cases, page thrashing may occur and impact performance. An introduction to CUDA Unified Memory is available on Parallel Forall.

This feature has the following limitations:

• Use of managed memory applies only to dynamically-allocated data. Static data (C static and extern variables, Fortran module, common block and save variables) and function local data is still handled by the OpenACC runtime. Dynamically allocated Fortran local variables and Fortran allocatable arrays are implicitly managed but Fortran array pointers are not.
• Given an allocatable aggregate with a member that points to local, global or static data, compiling with -⁠ta=tesla:managed and attempting to access memory through that pointer from the compute kernel will cause a failure at runtime.
• C++ virtual functions are not supported.
• The -⁠ta=tesla:managed compiler option must be used to compile the files in which variables are allocated, even if there is no OpenACC code in the file.

This feature has the following additional limitations when used with NVIDIA Kepler GPUs:

• Data motion on Kepler GPUs is achieved through fast pinned asynchronous data transfers; from the program's perspective, however, the transfers are synchronous.
• The PGI runtime enforces synchronous execution of kernels when -⁠ta=tesla:managed is used on a system with a Kepler GPU. This situation may result in slower performance because of the extra synchronizations and decreased overlap between CPU and GPU.
• The total amount of managed memory is limited to the amount of available device memory on Kepler GPUs.

This feature is not supported on NVIDIA Fermi GPUs.

CUDA Unified Memory Pool Allocator

Dynamic memory allocations are made using cudaMallocManaged(), a routine which has higher overhead than allocating non-unified memory using cudaMalloc(). The more calls to cudaMallocManaged(), the more significant the impact on performance.

To mitigate the overhead of cudaMallocManaged() calls, both -⁠ta=tesla:managed and -⁠ta=tesla:pinned use a CUDA Unified Memory pool allocator to minimize the number of calls to cudaMallocManaged(). The pool allocator is enabled by default. It can be disabled, or its behavior modified, using these environment variables:

Multicore Support

PGI Accelerator OpenACC compilers support the option -⁠ta=multicore, to set the target accelerator for OpenACC programs to the host multicore. This will compile OpenACC compute regions for parallel execution across the cores of the host processor or processors. The host multicore will be treated as a shared-memory accelerator, so the data clauses (copy, copyin, copyout, create) will be ignored and no data copies will be executed.

By default, -⁠ta=multicore will generate code that will use all the available cores of the processor. If the compute region specifies a value in the num_gangs clause, the minimum of the num_gangs value and the number of available cores will be used. At runtime, the number of cores can be limited by setting the environment variable ACC_NUM_CORES to a constant integer value. If an OpenACC compute construct appears lexically within an OpenMP parallel construct, the OpenACC compute region will generate sequential code. If an OpenACC compute region appears dynamically within an OpenMP region or another OpenACC compute region, the program may generate many more threads than there are cores, and may produce poor performance.

The ACC_BIND environment variable is set by default with -⁠ta=multicore; ACC_BIND has similiar behavior to MP_BIND for OpenMP.

The -⁠ta=multicore option differs from the -⁠ta=host option in that -⁠ta=host generates sequential code for the OpenACC compute regions.

Default Compute Capability

The default compute capability list for OpenACC and CUDA Fortran compilation for NVIDIA Tesla targets is cc30, cc35, and cc50. If CUDA 8.0 is specified using the cuda8.0 sub-option, cc60 is added to the default compute capability list. The generation of device code can be time consuming, so you may notice an increase in compile time. You can override the default by specifying one or more compute capabilities using either command-line options or an rcfile.

To change the default with a command-line option, provide a comma-separated list of compute capabilities to -⁠ta=tesla: for OpenACC or -⁠Mcuda= for CUDA Fortran.

To change the default with an rcfile, set the DEFCOMPUTECAP value to a blank-separated list of compute capabilities in the siterc file located in your installation's bin directory:

set DEFCOMPUTECAP=30 35 50;

Alternatively, if you don't have permissions to change the siterc file, you can add the DEFCOMPUTECAP definition to a separate .mypgirc file (mypgi_rc on Windows) in your home directory.

OpenACC 2.5

The PGI compilers implement OpenACC 2.5 as defined in The OpenACC Application Programming Interface, Version 2.5, August 2013, http://www.openacc.org, with the exception that these features are not yet supported:

• nested parallelism
• declare link
• enforcement of the new cache clause restriction that all references to listed variables must lie within the region being cached

Support for Profiler/Trace Tool Interface

PGI compilers support the OpenACC 2.5 version of the OpenACC Profiler/Trace Tools Interface. This is the interface used by the PGI profiler to collect performance measurements of OpenACC programs.

ACC routine directive limitations

• Fortran assumed-shape arguments are not yet supported.

Clause Support Limitations