mirror of
https://github.com/boostorg/math.git
synced 2026-01-19 04:22:09 +00:00
1324401c92
Add SYCL testing of ariy functions Add CUDA testing of airy functions Add NVRTC testing of airy functions Add GPU support to ellint rc Add GPU support to ellint rd Add GPU support to ellint rf Add GPU support to ellint rg Add GPU support to ellint rj Add GPU support to ellint d Add GPU support to ellint_1 Markup forward and add ellint_3 return type def for NVRTC platform Add CUDA testing of ellint 1 NVRTC fixes Add NVRTC testing of ellint_1 Add GPU support to ellint_2 Add CUDA testing of ellint_2 Fix NVRTC errors Add NVRTC testing of ellint_2 Add GPU support to atanh Add GPU support to ellint_3 Add NVRTC testing of ellint_3 Add CUDA testing of ellint_3 Replace use of static const char* Add SYCL testing of ellint_1 Add SYCL testing of ellint 2 with slight tolerance bump Remove recursion from ellint_rj Add ellint_d CUDA testing Add NVRTC testing of ellint_d Add SYCL testing of ellint_d Remove SYCL ellint_3 support Update docs Add GPU support to jacobi zeta Add CUDA testing of jacobi zeta Add NVRTC testing of jacobi zeta Add SYCL testing of jacobi zeta Add GPU support to heuman_lambda Add NVRTC testing of heuman lambda Add CUDA testing of heuman_lambda Add SYCL testing of heuman lambda Add markers to docs Add marker for CUDA only functions in the docs
68 lines
2.5 KiB
Plaintext
68 lines
2.5 KiB
Plaintext
[section:gpu Support for GPU programming in Boost.Math]
|
|
|
|
[h4 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.
|
|
|
|
[h4 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.
|
|
|
|
[h4 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.
|
|
|
|
[endsect] [/section:gpu Support for GPU programming in Boost.Math]
|
|
|
|
[/
|
|
Copyright 2024. Matt Borland
|
|
Distributed under the Boost Software License, Version 1.0.
|
|
(See accompanying file LICENSE_1_0.txt or copy at
|
|
http://www.boost.org/LICENSE_1_0.txt).
|
|
]
|
|
|