Boost C++ Libraries

GPU Support

Selected functions, distributions, tools, etc. support running on both host and devices. These functions will have the annotation BOOST_MATH_GPU_ENABLED or BOOST_MATH_CUDA_ENABLED next to their individual documentation. Functions marked with BOOST_MATH_GPU_ENABLED are tested using CUDA (both NVCC and NVRTC) as well as SYCL to provide a wide range of support. Functions marked with BOOST_MATH_CUDA_ENABLED are few, but due to its restrictions SYCL is unsupported.

Policies

The default policy on all devices is ignore error due to the lack of throwing ability. A user can specify their own policy like usual, but when the code is run on device it will be ignored.

How to build with device support

When compiling with CUDA or SYCL you will have to ensure that your code is being run inside of a kernel function. It is not enough to simply compile existing code with the NVCC compiler to run the code on the device. A simple CUDA kernel to run the Beta Distribution CDF on NVCC would be:

__global__ void cuda_beta_dist(const double* in, double* out, int num_elements)
{
    const int i = blockDim.x * blockIdx.x + threadIdx.x;

    if (i < num_elements)
    {
        out[i] = cdf(boost::math::beta_distribution<double>(), in[i]);
    }
}

And on CUDA on NVRTC:

const char* cuda_kernel = R"(
#include <boost/math/distributions/beta.hpp>
extern "C" __global__
void test_beta_dist_kernel(const double* in, double* out, int num_elements)
{
    const int i = blockDim.x * blockIdx.x + threadIdx.x;
    if (i < num_elements)
    {
        out[i] = boost::math::cdf(boost::math::beta_distribution<double>(), in[i]);
    }
}
)";

And lastly on SYCL:

void sycl_beta_dist(const double* in, double* out, int num_elements, sycl::queue& q)
{
    q.submit([&](sycl::handler& h) {
        h.parallel_for(sycl::range<1>(num_elements), [=](sycl::id<1> i) {
            out[i] = boost::math::cdf(boost::math::beta_distribution<double>(), in[i]);
        });
    });
}

Once your kernel function has been written then use the framework mechanism for launching the kernel.