GPU acceleration of C++ Parallel Algorithms is enabled with the -⁠stdpar=gpu command-line option to NVC++. If -⁠stdpar=gpu is specified (or -⁠stdpar without an argument), almost all algorithms that use a parallel execution policy are compiled for offloading to run in parallel on an NVIDIA GPU:

  nvc++ -stdpar=gpu program.cpp -o program 

  nvc++ -stdpar program.cpp -o program 

Acceleration of C++ Parallel Algorithms with multicore CPUs is enabled with the -⁠stdpar=multicore command-line option to NVC++. If -⁠stdpar=multicore is specified, all algorithms that use a parallel execution policy are compiled to run in on a multicore CPU:

  nvc++ -stdpar=multicore program.cpp -o program 

Functions to be executed on the GPU within parallel algorithms do not need any __device__ annotations or other special markings to be compiled for GPU execution. The NVC++ compiler walks the call graph for each source file and automatically infers which functions must be compiled for GPU execution.

However, this only works when the compiler can see the function definition in the same source file where the function is called. This is true for most inline functions and template functions but may fail when functions are defined in a different source file or linked in from an external library. You need to be aware of this when formulating parallel algorithms invocations that you expect to be offloaded and accelerated on NVIDIA GPUs.

NVC⁠+⁠+ relies on CUDA Unified Memory for all data movement between CPU and GPU memory. Through support in both the CUDA device driver and the NVIDIA GPU hardware, the CUDA Unified Memory manager automatically moves some types of data based on usage. Some restrictions on data accesses from the parallel algorithms are in place depending on the exact memory capability of the compilation target. This section details such restrictions, most of which are related to compilation with -⁠gpu=managed.

A general discussion of the memory models which are now supported can be found in the NVIDIA HPC Compiler User’s Guide, available online at docs.nvidia.com/hpc-sdk.

std::vector<int> v = ...;
std::sort(std::execution::par, 
          v.begin(), v.end()); // Okay, accesses heap memory.
      

std::array<int, 1024> a = ...;
std::sort(std::execution::par, 
           a.begin(), a.end()); // Fails on targets where only managed
                                // memory is unified, array is on a CPU
                                // stack inaccessible from GPU.
                                // Works correctly on targets where system 
                                // memory is unified. 
      

void saxpy(float* x, float* y, int N, float a) {
   std::transform(std::execution::par_unseq, x, x + N, y, y,
                  [&](float xi, float yi){ return a * xi + yi; });
}
      

void saxpy(float* x, float* y, int N, float a) {
   std::transform(std::execution::par_unseq, x, x + N, y, y,
                  [=](float xi, float yi){ return a * xi + yi; });
}
      

int globvar = 123;
void foo() {
  auto r = std::views::iota(0, N);
  std::for_each(std::execution::par_unseq, r.begin(), r.end(),
                [](auto i) {
                       if (i == N - 1)
                         globvar += 345;
          });
  // globvar is equal to 468.
}
      

Functions compiled to run on either the CPU or the GPU must be compiled into two different versions, one with the CPU machine instructions and one with the GPU machine instructions.

In the current implementation, a function pointer either points to the CPU or the GPU version of the functions. This causes problems if you attempt to pass function pointers between CPU and GPU code. You might inadvertently pass a pointer to the CPU version of the function to GPU code. In the future, it may be possible to automatically and seamlessly support the use of function pointers across CPU and GPU code boundaries, but it is not supported in the current implementation.

Function pointers can’t be passed to Parallel Algorithms to be run on the GPU, and functions may not be called through a function pointer within GPU code. For example, the following code example won’t work correctly:

  void square(int& x) { x = x * x; }
  void square_all(std::vector<int>& v) {
    std::for_each(std::execution::par_unseq,
                  v.begin(), v.end(), &square);
  }

It passes a pointer to the CPU version of the function square to a parallel for_each algorithm invocation. When the algorithm is parallelized and offloaded to the GPU, the program fails to resolve the function pointer to the GPU version of square.

You can often solve this issue by using a function object, which is an object with a function call operator. The function object's call operator is resolved at compile time to the GPU version of the function, instead of being resolved at run time to the incorrect CPU version of the function as in the previous example. For example, the following code example works:

  struct squared {
    void operator()(int& x) const { x = x * x; }
  };
  void square_all(std::vector<int>& v) {
    std::for_each(std::execution::par_unseq,
                  v.begin(), v.end(), squared{});
  }

Another possible workaround is to change the function to a lambda, because a lambda is implemented as a nameless function object:

  void square_all(std::vector<int>& v) {
    std::for_each(std::execution::par_unseq, v.begin(), v.end(),
                  [](int& x) { x = x * x; });
  }

If the function in question is too big to be converted to a function object or a lambda, then it should be possible to wrap the call to the function in a lambda:

  void compute(int& x) {
    // Assume lots and lots of code here.
  }
  void compute_all(std::vector<int>& v) {
    std::for_each(std::execution::par_unseq, v.begin(), v.end(),
                  [](int& x) { compute(x); });
  }

No function pointers are used in this example.

The restriction on calling a function through a function pointer unfortunately means passing polymorphic objects from CPU code to GPU-accelerated Parallel Algorithms is not currently supported, as virtual tables are implemented using function pointers.

The C++ Standard requires that the iterators passed to most C++ Parallel Algorithms be forward iterators. However, C++ Parallel Algorithms on GPUs only works with random access iterators. Passing a forward iterator or a bidirectional iterator to a GPU/CPU-accelerated Parallel Algorithm results in a compilation error. Passing raw pointers or Standard Library random access iterators to the algorithms has the best performance, but most other random-access iterators work correctly.

Large parts of the C++ Standard Library can be used with stdpar on GPUs.

• std::atomic<T> objects within GPU code work provided that T is a four-byte or eight-byte integer type.
• Math functions that operate on floating-point types—such as sin, cos, log, and most of the other functions declared in <cmath> —can be used in GPU code and resolve to the same implementations that are used in CUDA C++ programs.
• std::complex, std::tuple, std::pair, std::optional, std::variant, and <type_traits>, are supported and work as expected in GPU code.

The parts of the C++ Standard Library that aren’t supported in GPU code include I/O functions and in general any function that accesses the CPU operating system. As a special case, basic printf calls can be used within GPU code and leverage the same implementation that is used in NVIDIA CUDA C++.

As with most other GPU programming models, throwing and catching C++ exceptions is not supported within Parallel Algorithm invocations that are offloaded to the GPU.

Unlike some other GPU programming models where try/catch blocks and throw expressions are compilation errors, exception code does compile but with non-standard behavior. Catch clauses are ignored, and throw expressions abort the GPU kernel if actually executed. Exceptions in CPU code work without restrictions.

Multi-dimensional spans (std::mdspan) enable customizable multi-dimensional access to data. This feature was added to C++23 (see P0009 and follow-on papers). A Gentle Introduction to mdspan gives a tutorial. The reference mdspan implementation https://github.com/kokkos/mdspan also has many useful examples.

nvc++ provides an implementation available in the <experimental/mdspan> namespace that works with C++17 or newer. It enables applications that are not targeting the C++23 version of the standard to use mdspan.

nvc++ also provides the P0009R17 version of submdspan, which only works for the mdspan layouts in C++23; that is, it does not implement C++26 submdspan (P2630) yet.

C++23’s mdspan uses operator[] for array access. For example, if A is a rank-2 mdspan, and i and j are integers, then A[i, j] accesses the element of A at row i and column j. Before C++23, operator[] was only allowed to take one argument. C++23 changed the language to permit any number of arguments (zero or more). nvc++ does not support this new language feature. As a result, the implementation of mdspan provided by nvc++ permits use of operator() as a fall-back (e.g., A(i, j) instead of A[i, j]). Users may enable this fall-back manually, by defining the macro MDSPAN_USE_PAREN_OPERATOR to 1 before including any mdspan headers.

The following example (godbolt):

#include <experimental/mdspan>
#include <iostream>


namespace stdex = std::experimental;


int main() {
 std::array d{
   0, 5, 1,
   3, 8, 4,
   2, 7, 6,
 };


 stdex::mdspan m{d.data(), stdex::extents{3, 3}};
 static_assert(m.rank()==2, "Rank is two");


 for (std::size_t i = 0; i < m.extent(0); ++i)
   for (std::size_t j = 0; j < m.extent(1); ++j)
     std::cout << "m(" << i << ", " << j << ") == " << m(i, j) << "\n";


 return 0;
}

is compiled as follows

nvc++ -std=c++17 -o example example.cpp

and outputs

m(0, 0) == 0
m(0, 1) == 5
m(0, 2) == 1
m(1, 0) == 3
m(1, 1) == 8
m(1, 2) == 4
m(2, 0) == 2
m(2, 1) == 7
m(2, 2) == 6

P2300 - std::execution proposes a model of asynchronous programming for adoption into the C++26 Standard. For an introduction to this feature, see Design - user side section of the proposal. The NVIDIA implementation of Senders and receivers is open source and its repository contains many useful examples. nvc++ provides access to the NVIDIA implementation which works in C++20 or newer. Since the proposal is still evolving, our implementation is not stable. It is experimental in nature and will change to follow the proposal closely without any warning. The NVIDIA implementation is structured as follows:

The following example (godbolt) builds a task graph in which two different vectors, v0 and v1, are concurrently modified in bulk, using a CPU thread pool and a GPU stream context, respectively. This graph then transfers execution to the CPU thread pool, and adds both vectors into v2 on the CPU, returning the sum of all elements:

int main()
{
  // Declare a pool of 8 worker CPU threads:
  exec::static_thread_pool pool(8);


  // Declare a GPU stream context:
  nvexec::stream_context stream_ctx{};


  // Get a handle to the thread pool:
  auto cpu_sched = pool.get_scheduler();
  auto gpu_sched = stream_ctx.get_scheduler();


  // Declare three dynamic array with N elements
  std::size_t N = 5;
  std::vector<int> v0 {1, 1, 1, 1, 1};
  std::vector<int> v1 {2, 2, 2, 2, 2};
  std::vector<int> v2 {0, 0, 0, 0, 0};


  // Describe some work:
  auto work = stdexec::when_all(
    // Double v0 on the CPU
    stdexec::just()
      | exec::on(cpu_sched,
                 stdexec::bulk(N, [v0 = v0.data()](std::size_t i) {
                   v0[i] *= 2;
      })),
    // Triple v1 on the GPU
    stdexec::just()
      | exec::on(gpu_sched,
                 stdexec::bulk(N, [v1 = v1.data()](std::size_t i) {
                   v1[i] *= 3;
    }))
  )
  | stdexec::transfer(cpu_sched)
  // Add the two vectors into the output vector v2 = v0 + v1:
  | stdexec::bulk(N, [&](std::size_t i) { v2[i] = v0[i] + v1[i]; })
  | stdexec::then([&] {
    int r = 0;
    for (std::size_t i = 0; i < N; ++i) r += v2[i];
      return r;
  });
  auto [sum] = stdexec::sync_wait(work).value();
  // Print the results:
  std::printf("sum = %d\n", sum);
  for (int i = 0; i < N; ++i) {
    std::printf("v0[%d] = %d, v1[%d] = %d, v2[%d] = %d\n",
                i, v0[i], i, v1[i], i, v2[i]);
  }
  return 0;
}

is compiled as follows:

nvc++ --stdpar=gpu --experimental-stdpar -std=c++20 -o example example.cpp

and outputs:

sum = 40
v0[0] = 2, v1[0] = 6, v2[0] = 8
v0[1] = 2, v1[1] = 6, v2[1] = 8
v0[2] = 2, v1[2] = 6, v2[2] = 8
v0[3] = 2, v1[3] = 6, v2[3] = 8
v0[4] = 2, v1[4] = 6, v2[4] = 8

P1673 - A free function linear algebra interface based on the BLAS proposes standardizing an idiomatic C++ interface based on std::mdspan for a subset of the Basic Linear Algebra Subroutines (BLAS) standard. For an introduction to this feature, see P1673 (C++ linear algebra library) background & motivation. There are many useful examples available in $HPCSDK_HOME/examples/stdpar/stdblas and in the repository of the reference implementation. A detailed documentation is available at $HPCSDK_HOME/compilers/include/experimental/__p1673_bits/README.md. nvc++ provides access to the NVIDIA implementation which works in C++17 or newer. Since the proposal is still evolving, our implementation is not stable. It is experimental in nature and will change to follow the proposal closely without any warning. To use the linear algebra library facilities, a suitable linear algebra library must be linked: cuBLAS for GPU execution via the -cudalib=cublas flag, and a CPU BLAS library for CPU execution. The HPC SDK bundles OpenBLAS which may be linked using the -lblas linker flag.

The following example (godbolt):

#include <experimental/mdspan>
#include <experimental/linalg>
#include <vector>
#include <array>


namespace stdex = std::experimental;


int main()
{
 constexpr size_t N = 4;
 constexpr size_t M = 2;


 std::vector<double> A_vec(N*M);
 std::vector<double> x_vec(M);
 std::array<double, N> y_vec(N);


 stdex::mdspan A(A_vec.data(), N, M);
 stdex::mdspan x(x_vec.data(), M);
 stdex::mdspan y(y_vec.data(), N);


 for(int i = 0; i < A.extent(0); ++i)
   for(int j = 0; j < A.extent(1); ++j)
     A(i,j) = 100.0 * i + j;


 for(int j = 0; j < x.extent(0); ++j) x(j) = 1.0 * j;
 for(int i = 0; i < y.extent(0); ++i) y(i) = -1.0 * i;


 stdex::linalg::matrix_vector_product(A, x, y); // y = A * x


 // y = 0.5 * y + 2 * A * x
 stdex::linalg::matrix_vector_product(std::execution::par,
 stdex::linalg::scaled(2.0, A), x,
 stdex::linalg::scaled(0.5, y), y);


 // Print the results:
 for (int i = 0; i < N; ++i) std::printf("y[%d] = %f\n", i, y(i));
 return 0;
}

is compiled as follows for GPU execution:

nvc++ -std=c++17 -stdpar=gpu -cudalib=cublas -o example example.cpp

And as follows for CPU execution:

nvc++ -std=c++17 -stdpar=multicore -o example example.cpp -lblas

and produces the same outputs in both cases:

y[0] = 2.500000
y[1] = 252.500000
y[2] = 502.500000
y[3] = 752.500000

The NVC⁠+⁠+ compiler can automatically offload C⁠+⁠+ Parallel Algorithms to NVIDIA GPUs based on the Volta architecture or newer. These architectures include features -- such as independent thread scheduling and hardware optimizations for CUDA Unified Memory -- that were specifically designed to support high-performance, general-purpose parallel programming models like the C⁠+⁠+ Parallel Algorithms.

The NVC⁠+⁠+ compiler provides limited support for C⁠+⁠+ Parallel Algorithms on the Pascal architecture, which does not have the independent thread scheduling necessary to properly support the std::execution::par policy. When compiling for the Pascal architecture (-⁠gpu=cc60), NVC⁠+⁠+ compiles algorithms with the std::execution::par policy for serial execution on the CPU. Only algorithms with the std::execution::par_unseq policy will be scheduled to run on Pascal GPUs.

By default, NVC⁠+⁠+ auto-detects and generates GPU code for the type of GPU that is installed on the system on which the compiler is running. To generate code for a specific GPU architecture, which may be necessary when the application is compiled and run on different systems, add the -⁠gpu=ccXX command-line option.

The NVC++ compiler is built on CUDA libraries and technologies and uses CUDA to accelerate C++ Parallel Algorithms on NVIDIA GPUs. A GPU-accelerated system on which NVC++-compiled applications are to be run must have a CUDA 11.2 or newer device driver installed.

The NVIDIA HPC SDK compilers ship with an integrated CUDA toolchain, header files, and libraries to use during compilation, so it is not necessary to have a CUDA Toolkit installed on the system.

When -⁠stdpar is specified, NVC++ compiles using the CUDA toolchain version that best matches the CUDA driver installed on the system on which compilation is performed. To compile using a different version of the CUDA toolchain, use the -⁠gpu=cudaX.Y option. For example, use the -⁠gpu=cuda11.8 option to specify that your program should be compiled for a CUDA 11.8 system using the CUDA 11.8 toolchain.