mirror of
https://github.com/boostorg/math.git
synced 2026-01-19 04:22:09 +00:00
122 lines
11 KiB
HTML
122 lines
11 KiB
HTML
<html>
|
||
<head>
|
||
<meta http-equiv="Content-Type" content="text/html; charset=UTF-8">
|
||
<title>Support for GPU programming in Boost.Math</title>
|
||
<link rel="stylesheet" href="../math.css" type="text/css">
|
||
<meta name="generator" content="DocBook XSL Stylesheets Vsnapshot">
|
||
<link rel="home" href="../index.html" title="Math Toolkit 4.2.1">
|
||
<link rel="up" href="../overview.html" title="Chapter 1. Overview">
|
||
<link rel="prev" href="error_handling.html" title="Error Handling">
|
||
<link rel="next" href="compilers_overview.html" title="Compilers">
|
||
<meta name="viewport" content="width=device-width, initial-scale=1">
|
||
</head>
|
||
<body bgcolor="white" text="black" link="#0000FF" vlink="#840084" alink="#0000FF">
|
||
<table cellpadding="2" width="100%"><tr>
|
||
<td valign="top"><img alt="Boost C++ Libraries" width="277" height="86" src="../../../../../boost.png"></td>
|
||
<td align="center"><a href="../../../../../index.html">Home</a></td>
|
||
<td align="center"><a href="../../../../../libs/libraries.htm">Libraries</a></td>
|
||
<td align="center"><a href="http://www.boost.org/users/people.html">People</a></td>
|
||
<td align="center"><a href="http://www.boost.org/users/faq.html">FAQ</a></td>
|
||
<td align="center"><a href="../../../../../more/index.htm">More</a></td>
|
||
</tr></table>
|
||
<hr>
|
||
<div class="spirit-nav">
|
||
<a accesskey="p" href="error_handling.html"><img src="../../../../../doc/src/images/prev.png" alt="Prev"></a><a accesskey="u" href="../overview.html"><img src="../../../../../doc/src/images/up.png" alt="Up"></a><a accesskey="h" href="../index.html"><img src="../../../../../doc/src/images/home.png" alt="Home"></a><a accesskey="n" href="compilers_overview.html"><img src="../../../../../doc/src/images/next.png" alt="Next"></a>
|
||
</div>
|
||
<div class="section">
|
||
<div class="titlepage"><div><div><h2 class="title" style="clear: both">
|
||
<a name="math_toolkit.gpu"></a><a class="link" href="gpu.html" title="Support for GPU programming in Boost.Math">Support for GPU programming in Boost.Math</a>
|
||
</h2></div></div></div>
|
||
<h5>
|
||
<a name="math_toolkit.gpu.h0"></a>
|
||
<span class="phrase"><a name="math_toolkit.gpu.gpu_support"></a></span><a class="link" href="gpu.html#math_toolkit.gpu.gpu_support">GPU
|
||
Support</a>
|
||
</h5>
|
||
<p>
|
||
Selected functions, distributions, tools, etc. support running on both host
|
||
and devices. These functions will have the annotation <code class="computeroutput"><span class="identifier">BOOST_MATH_GPU_ENABLED</span></code>
|
||
or <code class="computeroutput"><span class="identifier">BOOST_MATH_CUDA_ENABLED</span></code>
|
||
next to their individual documentation. Functions marked with <code class="computeroutput"><span class="identifier">BOOST_MATH_GPU_ENABLED</span></code> are tested using CUDA
|
||
(both NVCC and NVRTC) as well as SYCL to provide a wide range of support. Functions
|
||
marked with <code class="computeroutput"><span class="identifier">BOOST_MATH_CUDA_ENABLED</span></code>
|
||
are few, but due to its restrictions SYCL is unsupported.
|
||
</p>
|
||
<h5>
|
||
<a name="math_toolkit.gpu.h1"></a>
|
||
<span class="phrase"><a name="math_toolkit.gpu.policies"></a></span><a class="link" href="gpu.html#math_toolkit.gpu.policies">Policies</a>
|
||
</h5>
|
||
<p>
|
||
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.
|
||
</p>
|
||
<h5>
|
||
<a name="math_toolkit.gpu.h2"></a>
|
||
<span class="phrase"><a name="math_toolkit.gpu.how_to_build_with_device_support"></a></span><a class="link" href="gpu.html#math_toolkit.gpu.how_to_build_with_device_support">How
|
||
to build with device support</a>
|
||
</h5>
|
||
<p>
|
||
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:
|
||
</p>
|
||
<pre class="programlisting"><span class="identifier">__global__</span> <span class="keyword">void</span> <span class="identifier">cuda_beta_dist</span><span class="special">(</span><span class="keyword">const</span> <span class="keyword">double</span><span class="special">*</span> <span class="identifier">in</span><span class="special">,</span> <span class="keyword">double</span><span class="special">*</span> <span class="identifier">out</span><span class="special">,</span> <span class="keyword">int</span> <span class="identifier">num_elements</span><span class="special">)</span>
|
||
<span class="special">{</span>
|
||
<span class="keyword">const</span> <span class="keyword">int</span> <span class="identifier">i</span> <span class="special">=</span> <span class="identifier">blockDim</span><span class="special">.</span><span class="identifier">x</span> <span class="special">*</span> <span class="identifier">blockIdx</span><span class="special">.</span><span class="identifier">x</span> <span class="special">+</span> <span class="identifier">threadIdx</span><span class="special">.</span><span class="identifier">x</span><span class="special">;</span>
|
||
|
||
<span class="keyword">if</span> <span class="special">(</span><span class="identifier">i</span> <span class="special"><</span> <span class="identifier">num_elements</span><span class="special">)</span>
|
||
<span class="special">{</span>
|
||
<span class="identifier">out</span><span class="special">[</span><span class="identifier">i</span><span class="special">]</span> <span class="special">=</span> <span class="identifier">cdf</span><span class="special">(</span><span class="identifier">boost</span><span class="special">::</span><span class="identifier">math</span><span class="special">::</span><span class="identifier">beta_distribution</span><span class="special"><</span><span class="keyword">double</span><span class="special">>(),</span> <span class="identifier">in</span><span class="special">[</span><span class="identifier">i</span><span class="special">]);</span>
|
||
<span class="special">}</span>
|
||
<span class="special">}</span>
|
||
</pre>
|
||
<p>
|
||
And on CUDA on NVRTC:
|
||
</p>
|
||
<pre class="programlisting"><span class="keyword">const</span> <span class="keyword">char</span><span class="special">*</span> <span class="identifier">cuda_kernel</span> <span class="special">=</span> <span class="identifier">R</span><span class="string">"(
|
||
#include <boost/math/distributions/beta.hpp>
|
||
extern "</span><span class="identifier">C</span><span class="string">" __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]);
|
||
}
|
||
}
|
||
)"</span><span class="special">;</span>
|
||
</pre>
|
||
<p>
|
||
And lastly on SYCL:
|
||
</p>
|
||
<pre class="programlisting"><span class="keyword">void</span> <span class="identifier">sycl_beta_dist</span><span class="special">(</span><span class="keyword">const</span> <span class="keyword">double</span><span class="special">*</span> <span class="identifier">in</span><span class="special">,</span> <span class="keyword">double</span><span class="special">*</span> <span class="identifier">out</span><span class="special">,</span> <span class="keyword">int</span> <span class="identifier">num_elements</span><span class="special">,</span> <span class="identifier">sycl</span><span class="special">::</span><span class="identifier">queue</span><span class="special">&</span> <span class="identifier">q</span><span class="special">)</span>
|
||
<span class="special">{</span>
|
||
<span class="identifier">q</span><span class="special">.</span><span class="identifier">submit</span><span class="special">([&](</span><span class="identifier">sycl</span><span class="special">::</span><span class="identifier">handler</span><span class="special">&</span> <span class="identifier">h</span><span class="special">)</span> <span class="special">{</span>
|
||
<span class="identifier">h</span><span class="special">.</span><span class="identifier">parallel_for</span><span class="special">(</span><span class="identifier">sycl</span><span class="special">::</span><span class="identifier">range</span><span class="special"><</span><span class="number">1</span><span class="special">>(</span><span class="identifier">num_elements</span><span class="special">),</span> <span class="special">[=](</span><span class="identifier">sycl</span><span class="special">::</span><span class="identifier">id</span><span class="special"><</span><span class="number">1</span><span class="special">></span> <span class="identifier">i</span><span class="special">)</span> <span class="special">{</span>
|
||
<span class="identifier">out</span><span class="special">[</span><span class="identifier">i</span><span class="special">]</span> <span class="special">=</span> <span class="identifier">boost</span><span class="special">::</span><span class="identifier">math</span><span class="special">::</span><span class="identifier">cdf</span><span class="special">(</span><span class="identifier">boost</span><span class="special">::</span><span class="identifier">math</span><span class="special">::</span><span class="identifier">beta_distribution</span><span class="special"><</span><span class="keyword">double</span><span class="special">>(),</span> <span class="identifier">in</span><span class="special">[</span><span class="identifier">i</span><span class="special">]);</span>
|
||
<span class="special">});</span>
|
||
<span class="special">});</span>
|
||
<span class="special">}</span>
|
||
</pre>
|
||
<p>
|
||
Once your kernel function has been written then use the framework mechanism
|
||
for launching the kernel.
|
||
</p>
|
||
</div>
|
||
<div class="copyright-footer">Copyright © 2006-2021 Nikhar Agrawal, Anton Bikineev, Matthew Borland,
|
||
Paul A. Bristow, Marco Guazzone, Christopher Kormanyos, Hubert Holin, Bruno
|
||
Lalande, John Maddock, Evan Miller, Jeremy Murphy, Matthew Pulver, Johan Råde,
|
||
Gautam Sewani, Benjamin Sobotta, Nicholas Thompson, Thijs van den Berg, Daryle
|
||
Walker and Xiaogang Zhang<p>
|
||
Distributed under the Boost Software License, Version 1.0. (See accompanying
|
||
file LICENSE_1_0.txt or copy at <a href="http://www.boost.org/LICENSE_1_0.txt" target="_top">http://www.boost.org/LICENSE_1_0.txt</a>)
|
||
</p>
|
||
</div>
|
||
<hr>
|
||
<div class="spirit-nav">
|
||
<a accesskey="p" href="error_handling.html"><img src="../../../../../doc/src/images/prev.png" alt="Prev"></a><a accesskey="u" href="../overview.html"><img src="../../../../../doc/src/images/up.png" alt="Up"></a><a accesskey="h" href="../index.html"><img src="../../../../../doc/src/images/home.png" alt="Home"></a><a accesskey="n" href="compilers_overview.html"><img src="../../../../../doc/src/images/next.png" alt="Next"></a>
|
||
</div>
|
||
</body>
|
||
</html>
|