<!DOCTYPE html
PUBLIC "-//W3C//DTD XHTML 1.0 Transitional//EN" "http://www.w3.org/TR/xhtml1/DTD/xhtml1-transitional.dtd">
<html xmlns="http://www.w3.org/1999/xhtml" lang="en-us" xml:lang="en-us">
<head>
<meta http-equiv="Content-Type" content="text/html; charset=utf-8"></meta>
<meta http-equiv="X-UA-Compatible" content="IE=edge"></meta>
<meta name="copyright" content="(C) Copyright 2005"></meta>
<meta name="DC.rights.owner" content="(C) Copyright 2005"></meta>
<meta name="DC.Type" content="concept"></meta>
<meta name="DC.Title" content="Thrust"></meta>
<meta name="abstract" content="The API reference guide for Thrust, the CUDA C++ template library."></meta>
<meta name="description" content="The API reference guide for Thrust, the CUDA C++ template library."></meta>
<meta name="DC.Coverage" content="CUDA API References"></meta>
<meta name="DC.subject" content="CUDA Thrust, CUDA Thrust installation, CUDA Thrust vectors, CUDA Thrust transforms, CUDA Thrust reductions, CUDA Thrust prefix-sums, CUDA Thrust reordering, CUDA Thrust fancy iterators, CUDA Thrust sorting, CUDA Thrust iterators"></meta>
<meta name="keywords" content="CUDA Thrust, CUDA Thrust installation, CUDA Thrust vectors, CUDA Thrust transforms, CUDA Thrust reductions, CUDA Thrust prefix-sums, CUDA Thrust reordering, CUDA Thrust fancy iterators, CUDA Thrust sorting, CUDA Thrust iterators"></meta>
<meta name="DC.Format" content="XHTML"></meta>
<meta name="DC.Identifier" content="abstract"></meta>
<link rel="stylesheet" type="text/css" href="../common/formatting/commonltr.css"></link>
<link rel="stylesheet" type="text/css" href="../common/formatting/site.css"></link>
<title>Thrust :: CUDA Toolkit Documentation</title>
<!--[if lt IE 9]>
<script src="../common/formatting/html5shiv-printshiv.min.js"></script>
<![endif]-->
<script type="text/javascript" charset="utf-8" src="../common/scripts/tynt/tynt.js"></script>
<script type="text/javascript" charset="utf-8" src="../common/formatting/jquery.min.js"></script>
<script type="text/javascript" charset="utf-8" src="../common/formatting/jquery.ba-hashchange.min.js"></script>
<script type="text/javascript" charset="utf-8" src="../common/formatting/jquery.scrollintoview.min.js"></script>
<script type="text/javascript" src="../search/htmlFileList.js"></script>
<script type="text/javascript" src="../search/htmlFileInfoList.js"></script>
<script type="text/javascript" src="../search/nwSearchFnt.min.js"></script>
<script type="text/javascript" src="../search/stemmers/en_stemmer.min.js"></script>
<script type="text/javascript" src="../search/index-1.js"></script>
<script type="text/javascript" src="../search/index-2.js"></script>
<script type="text/javascript" src="../search/index-3.js"></script>
<link rel="canonical" href="http://docs.nvidia.com/cuda/thrust/index.html"></link>
<link rel="stylesheet" type="text/css" href="../common/formatting/qwcode.highlight.css"></link>
</head>
<body>
<header id="header"><span id="company">NVIDIA</span><span id="site-title">CUDA Toolkit Documentation</span><form id="search" method="get" action="search">
<input type="text" name="search-text"></input><fieldset id="search-location">
<legend>Search In:</legend>
<label><input type="radio" name="search-type" value="site"></input>Entire Site</label>
<label><input type="radio" name="search-type" value="document"></input>Just This Document</label></fieldset>
<button type="reset">clear search</button>
<button id="submit" type="submit">search</button></form>
</header>
<div id="site-content">
<nav id="site-nav">
<div class="category closed"><a href="../index.html" title="The root of the site.">CUDA Toolkit
v6.5</a></div>
<div class="category"><a href="index.html" title="Thrust">Thrust</a></div>
<ul>
<li>
<div class="section-link"><a href="#introduction">1. Introduction</a></div>
<ul>
<li>
<div class="section-link"><a href="#installation-and-versioning">1.1. Installation and Versioning</a></div>
</li>
</ul>
</li>
<li>
<div class="section-link"><a href="#vectors">2. Vectors</a></div>
<ul>
<li>
<div class="section-link"><a href="#thrust-namespace">2.1. Thrust Namespace</a></div>
</li>
<li>
<div class="section-link"><a href="#iterators-and-static-dispatching">2.2. Iterators and Static Dispatching</a></div>
</li>
</ul>
</li>
<li>
<div class="section-link"><a href="#algorithms">3. Algorithms</a></div>
<ul>
<li>
<div class="section-link"><a href="#transformations">3.1. Transformations</a></div>
</li>
<li>
<div class="section-link"><a href="#reductions">3.2. Reductions</a></div>
</li>
<li>
<div class="section-link"><a href="#prefix-sums">3.3. Prefix-Sums</a></div>
</li>
<li>
<div class="section-link"><a href="#reordering">3.4. Reordering</a></div>
</li>
<li>
<div class="section-link"><a href="#sorting">3.5. Sorting</a></div>
</li>
</ul>
</li>
<li>
<div class="section-link"><a href="#fancy-iterators">4. Fancy Iterators</a></div>
<ul>
<li>
<div class="section-link"><a href="#constant-iterator">4.1. constant_iterator</a></div>
</li>
<li>
<div class="section-link"><a href="#counting-iterator">4.2. counting_iterator</a></div>
</li>
<li>
<div class="section-link"><a href="#transform-iterator">4.3. transform_iterator</a></div>
</li>
<li>
<div class="section-link"><a href="#permutation-iterator">4.4. permutation_iterator</a></div>
</li>
<li>
<div class="section-link"><a href="#zip-iterator">4.5. zip_iterator</a></div>
</li>
</ul>
</li>
<li>
<div class="section-link"><a href="#additional-resources">5. Additional Resources</a></div>
</li>
</ul>
</nav>
<div id="resize-nav"></div>
<nav id="search-results">
<h2>Search Results</h2>
<ol></ol>
</nav>
<div id="contents-container">
<div id="breadcrumbs-container">
<div id="eqn-warning">This document includes math equations
(highlighted in red) which are best viewed with <a target="_blank" href="https://www.mozilla.org/firefox">Firefox</a> version 4.0
or higher, or another <a target="_blank" href="http://www.w3.org/Math/Software/mathml_software_cat_browsers.html">MathML-aware
browser</a>. There is also a <a href="../../pdf/Thrust_Quick_Start_Guide.pdf">PDF version of this document</a>.
</div>
<div id="release-info">Thrust
(<a href="../../pdf/Thrust_Quick_Start_Guide.pdf">PDF</a>)
-
v6.5
(<a href="https://developer.nvidia.com/cuda-toolkit-archive">older</a>)
-
Last updated August 1, 2014
-
<a href="mailto:cudatools@nvidia.com?subject=CUDA Toolkit Documentation Feedback: Thrust">Send Feedback</a>
-
<span class="st_facebook"></span><span class="st_twitter"></span><span class="st_linkedin"></span><span class="st_reddit"></span><span class="st_slashdot"></span><span class="st_tumblr"></span><span class="st_sharethis"></span></div>
</div>
<article id="contents">
<div class="topic nested0" id="abstract"><a name="abstract" shape="rect">
<!-- --></a><h2 class="title topictitle1"><a href="#abstract" name="abstract" shape="rect">Thrust</a></h2>
<div class="body conbody"></div>
</div>
<div class="topic concept nested0" xml:lang="en-us" id="introduction"><a name="introduction" shape="rect">
<!-- --></a><h2 class="title topictitle1"><a href="#introduction" name="introduction" shape="rect">1. Introduction</a></h2>
<div class="body conbody">
<p class="p">Thrust is a C++ template library for CUDA based on the Standard Template Library (STL). Thrust allows you to implement high
performance parallel applications with minimal programming effort through a high-level interface that is fully interoperable
with CUDA C.
</p>
<p class="p">Thrust provides a rich collection of data parallel primitives such as scan, sort, and reduce, which can be composed together
to implement complex algorithms with concise, readable source code. By describing your computation in terms of these high-level
abstractions you provide Thrust with the freedom to select the most efficient implementation automatically. As a result, Thrust
can be utilized in rapid prototyping of CUDA applications, where programmer productivity matters most, as well as in production,
where robustness and absolute performance are crucial.
</p>
<p class="p">This document describes how to develop CUDA applications with Thrust. The tutorial is intended to be accessible, even if you
have limited C++ or CUDA experience.
</p>
</div>
<div class="topic concept nested1" xml:lang="en-us" id="installation-and-versioning"><a name="installation-and-versioning" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#installation-and-versioning" name="installation-and-versioning" shape="rect">1.1. Installation and Versioning</a></h3>
<div class="body conbody">
<p class="p">Installing the CUDA Toolkit will copy Thrust header files to the standard CUDA include directory for your system. Since Thrust
is a template library of header files, no further installation is necessary to start using Thrust.
</p>
<p class="p">In addition, new versions of Thrust continue to be available online through the GitHub <a class="xref" href="http://thrust.github.com/" target="_blank" shape="rect">Thrust project page</a>. The version of Thrust included in this version of the CUDA Toolkit corresponds to version 1.7.0 from the Thrust project
page.
</p>
</div>
</div>
</div>
<div class="topic concept nested0" xml:lang="en-us" id="vectors"><a name="vectors" shape="rect">
<!-- --></a><h2 class="title topictitle1"><a href="#vectors" name="vectors" shape="rect">2. Vectors</a></h2>
<div class="body conbody">
<p class="p">Thrust provides two <a class="xref" href="http://en.wikipedia.org/wiki/Vector_(C++)" target="_blank" shape="rect">vector</a> containers, <samp class="ph codeph">host_vector</samp> and <samp class="ph codeph">device_vector</samp>. As the names suggest, <samp class="ph codeph">host_vector</samp> is stored in host memory while <samp class="ph codeph">device_vector</samp> lives in GPU device memory. Thrust’s vector containers are just like <samp class="ph codeph">std::vector</samp> in the <a class="xref" href="http://www.sgi.com/tech/stl/Vector.html" target="_blank" shape="rect">C++ STL</a>. Like <samp class="ph codeph">std::vector</samp>, <samp class="ph codeph">host_vector</samp> and <samp class="ph codeph">device_vector</samp> are generic containers (able to store any data type) that can be resized dynamically. The following source code illustrates
the use of Thrust’s vector containers.
</p><pre xml:space="preserve">
#include <thrust/host_vector.h>
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-directive">#include <thrust/device_vector.h></span>
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-directive">#include <iostream></span>
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span> main(<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">void</span>)
{
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// H has storage for 4 integers</span>
thrust::host_vector<<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span>> H(4);
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// initialize individual elements</span>
H[0] = 14;
H[1] = 20;
H[2] = 38;
H[3] = 46;
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// H.size() returns the size of vector H</span>
std::cout << <span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-string">"H has size "</span> << H.size() << std::endl;
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// print contents of H</span>
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">for</span>(<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span> i = 0; i < H.size(); i++)
std::cout << <span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-string">"H["</span> << i << <span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-string">"] = "</span> << H[i] << std::endl;
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// resize H</span>
H.resize(2);
std::cout << <span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-string">"H now has size "</span> << H.size() << std::endl;
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// Copy host_vector H to device_vector D</span>
thrust::device_vector<<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span>> D = H;
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// elements of D can be modified</span>
D[0] = 99;
D[1] = 88;
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// print contents of D</span>
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">for</span>(<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span> i = 0; i < D.size(); i++)
std::cout << <span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-string">"D["</span> << i << <span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-string">"] = "</span> << D[i] << std::endl;
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// H and D are automatically deleted when the function returns</span>
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">return</span> 0;
}
</pre><p class="p">As this example shows, the <samp class="ph codeph">=</samp> operator can be used to copy a <samp class="ph codeph">host_vector</samp> to a <samp class="ph codeph">device_vector</samp> (or vice-versa). The <samp class="ph codeph">=</samp> operator can also be used to copy <samp class="ph codeph">host_vector</samp> to <samp class="ph codeph">host_vector</samp> or <samp class="ph codeph">device_vector</samp> to <samp class="ph codeph">device_vector</samp>. Also note that individual elements of a <samp class="ph codeph">device_vector</samp> can be accessed using the standard bracket notation. However, because each of these accesses requires a call to <samp class="ph codeph">cudaMemcpy</samp>, they should be used sparingly. We’ll look at some more efficient techniques later.
</p>
<p class="p">It’s often useful to initialize all the elements of a vector to a specific value, or to copy only a certain set of values
from one vector to another. Thrust provides a few ways to do these kinds of operations.
</p><pre xml:space="preserve">
#include <thrust/host_vector.h>
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-directive">#include <thrust/device_vector.h></span>
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-directive">#include <thrust/copy.h></span>
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-directive">#include <thrust/fill.h></span>
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-directive">#include <thrust/sequence.h></span>
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-directive">#include <iostream></span>
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span> main(<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">void</span>)
{
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// initialize all ten integers of a device_vector to 1</span>
thrust::device_vector<<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span>> D(10, 1);
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// set the first seven elements of a vector to 9</span>
thrust::fill(D.begin(), D.begin() + 7, 9);
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// initialize a host_vector with the first five elements of D</span>
thrust::host_vector<<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span>> H(D.begin(), D.begin() + 5);
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// set the elements of H to 0, 1, 2, 3, ...</span>
thrust::sequence(H.begin(), H.end());
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// copy all of H back to the beginning of D</span>
thrust::copy(H.begin(), H.end(), D.begin());
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// print D</span>
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">for</span>(<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span> i = 0; i < D.size(); i++)
std::cout << <span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-string">"D["</span> << i << <span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-string">"] = "</span> << D[i] << std::endl;
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">return</span> 0;
}
</pre><p class="p">Here we’ve illustrated use of the <samp class="ph codeph">fill</samp>, <samp class="ph codeph">copy</samp>, and <samp class="ph codeph">sequence</samp> functions. The <samp class="ph codeph">copy</samp> function can be used to copy a range of host or device elements to another host or device vector. Like the corresponding
STL function, <samp class="ph codeph">thrust::fill</samp> simply sets a range of elements to a specific value. Thrust’s <samp class="ph codeph">sequence</samp> function can be used to a create a sequence of equally spaced values.
</p>
</div>
<div class="topic concept nested1" xml:lang="en-us" id="thrust-namespace"><a name="thrust-namespace" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#thrust-namespace" name="thrust-namespace" shape="rect">2.1. Thrust Namespace</a></h3>
<div class="body conbody">
<p class="p">You’ll notice that we use things like <samp class="ph codeph">thrust::host_vector</samp> or <samp class="ph codeph">thrust::copy</samp> in our examples. The <samp class="ph codeph">thrust::</samp> part tells the C++ compiler that we want to look inside the <samp class="ph codeph">thrust </samp><dfn class="term">namespace</dfn> for a specific function or class. <a class="xref" href="http://en.wikipedia.org/wiki/Namespace_%28computer_science%29" target="_blank" shape="rect">Namespaces</a> are a nice way to avoid name collisions. For instance, <samp class="ph codeph">thrust::copy</samp> is different from <samp class="ph codeph">std::copy</samp> provided in the STL. C++ namespaces allow us to distinguish between these two <samp class="ph codeph">copy</samp> functions.
</p>
</div>
</div>
<div class="topic concept nested1" xml:lang="en-us" id="iterators-and-static-dispatching"><a name="iterators-and-static-dispatching" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#iterators-and-static-dispatching" name="iterators-and-static-dispatching" shape="rect">2.2. Iterators and Static Dispatching</a></h3>
<div class="body conbody">
<p class="p">In this section we used expressions like <samp class="ph codeph">H.begin()</samp> and <samp class="ph codeph">H.end()</samp> or offsets like <samp class="ph codeph">D.begin() + 7</samp>. The result of <samp class="ph codeph">begin()</samp> and <samp class="ph codeph">end()</samp> is called an <dfn class="term">iterator</dfn> in C++. In the case of vector containers, which are really just arrays, iterators can be thought of as pointers to array
elements. Therefore, <samp class="ph codeph">H.begin()</samp> is an iterator that points to the first element of the array stored inside the <samp class="ph codeph">H</samp> vector. Similarly, <samp class="ph codeph">H.end()</samp> points to the element <dfn class="term">one past</dfn> the last element of the <samp class="ph codeph">H</samp> vector.
</p>
<p class="p">Although vector iterators are similar to pointers they carry more information with them. Notice that we did not have to tell
<samp class="ph codeph">thrust::fill</samp> that it was operating on a <samp class="ph codeph">device_vector</samp> iterator. This information is captured in the type of the iterator returned by <samp class="ph codeph">D.begin()</samp> which is different than the type returned by <samp class="ph codeph">H.begin()</samp>. When a Thrust function is called, it inspects the type of the iterator to determine whether to use a host or a device implementation.
This process is known as <dfn class="term">static dispatching</dfn> since the host/device dispatch is resolved at compile time. Note that this implies that there is <dfn class="term">no runtime overhead</dfn> to the dispatch process.
</p>
<p class="p">You may wonder what happens when a “raw” pointer is used as an argument to a Thrust function. Like the STL, Thrust permits
this usage and it will dispatch the host path of the algorithm. If the pointer in question is in fact a pointer to device
memory then you’ll need to wrap it with <samp class="ph codeph">thrust::device_ptr</samp> before calling the function. For example:
</p><pre xml:space="preserve">
size_t N = 10;
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// raw pointer to device memory</span>
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span> * raw_ptr;
cudaMalloc((<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">void</span> **) &raw_ptr, N * <span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">sizeof</span>(<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span>));
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// wrap raw pointer with a device_ptr </span>
thrust::device_ptr<<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span>> dev_ptr(raw_ptr);
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// use device_ptr in thrust algorithms</span>
thrust::fill(dev_ptr, dev_ptr + N, (<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span>) 0);
</pre><p class="p">To extract a raw pointer from a <samp class="ph codeph">device_ptr</samp> the <samp class="ph codeph">raw_pointer_cast</samp> should be applied as follows:
</p><pre xml:space="preserve">
size_t N = 10;
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// create a device_ptr </span>
thrust::device_ptr<<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span>> dev_ptr = thrust::device_malloc<<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span>>(N);
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// extract raw pointer from device_ptr</span>
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span> * raw_ptr = thrust::raw_pointer_cast(dev_ptr);
</pre><p class="p">Another reason to distinguish between iterators and pointers is that iterators can be used to traverse many kinds of data
structures. For example, the STL provides a linked list container (<samp class="ph codeph">std::list</samp>) that provides bidirectional (but not random access) iterators. Although Thrust does not provide device implementations of
such containers, it is compatible with them.
</p><pre xml:space="preserve">
#include <thrust/device_vector.h>
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-directive">#include <thrust/copy.h></span>
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-directive">#include <list></span>
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-directive">#include <vector></span>
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span> main(<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">void</span>)
{
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// create an STL list with 4 values</span>
std::list<<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span>> stl_list;
stl_list.push_back(10);
stl_list.push_back(20);
stl_list.push_back(30);
stl_list.push_back(40);
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// initialize a device_vector with the list</span>
thrust::device_vector<<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span>> D(stl_list.begin(), stl_list.end());
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// copy a device_vector into an STL vector</span>
std::vector<<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span>> stl_vector(D.size());
thrust::copy(D.begin(), D.end(), stl_vector.begin());
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">return</span> 0;
}
</pre><p class="p"><dfn class="term">For Future Reference:</dfn> The iterators we’ve covered so far are useful, but fairly basic. In addition to these normal iterators, Thrust also provides
a collection of <dfn class="term">fancy iterators</dfn> with names like <samp class="ph codeph">counting_iterator</samp> and <samp class="ph codeph">zip_iterator</samp>. While they look and feel like normal iterators, fancy iterators are capable of more exciting things. We’ll revisit this
topic later in the tutorial.
</p>
</div>
</div>
</div>
<div class="topic concept nested0" xml:lang="en-us" id="algorithms"><a name="algorithms" shape="rect">
<!-- --></a><h2 class="title topictitle1"><a href="#algorithms" name="algorithms" shape="rect">3. Algorithms</a></h2>
<div class="body conbody">
<p class="p">Thrust provides a large number of common parallel algorithms. Many of these algorithms have direct analogs in the STL, and
when an equivalent STL function exists, we choose the name (e.g. <samp class="ph codeph">thrust::sort</samp> and <samp class="ph codeph">std::sort</samp>).
</p>
<p class="p">All algorithms in Thrust have implementations for both host and device. Specifically, when a Thrust algorithm is invoked with
a host iterator, then the host path is dispatched. Similarly, a device implementation is called when a device iterator is
used to define a range.
</p>
<p class="p">With the exception of <samp class="ph codeph">thrust::copy</samp>, which can copy data between host and device, all iterator arguments to a Thrust algorithm should live in the same place:
either all on the host or all on the device. When this requirement is violated the compiler will produce an error message.
</p>
</div>
<div class="topic concept nested1" xml:lang="en-us" id="transformations"><a name="transformations" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#transformations" name="transformations" shape="rect">3.1. Transformations</a></h3>
<div class="body conbody">
<p class="p">Transformations are algorithms that apply an operation to each element in a set of (zero or more) input ranges and then stores
the result in a destination range. One example we have already seen is <samp class="ph codeph">thrust::fill</samp>, which sets all elements of a range to a specified value. Other transformations include <samp class="ph codeph">thrust::sequence</samp>, <samp class="ph codeph">thrust::replace</samp>, and of course <samp class="ph codeph">thrust::transform</samp>. Refer to the <a class="xref" href="http://thrust.github.com/doc/group__transformations.html" target="_blank" shape="rect">documentation</a> for a complete listing.
</p>
<p class="p">The following source code demonstrates several of the transformation algorithms. Note that <samp class="ph codeph">thrust::negate</samp> and <samp class="ph codeph">thrust::modulus</samp> are known as functors in C++ terminology. Thrust provides these and other common functors like <samp class="ph codeph">plus</samp> and <samp class="ph codeph">multiplies</samp> in the file <samp class="ph codeph">thrust/functional.h</samp>.
</p><pre xml:space="preserve">
#include <thrust/device_vector.h>
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-directive">#include <thrust/transform.h></span>
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-directive">#include <thrust/sequence.h></span>
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-directive">#include <thrust/copy.h></span>
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-directive">#include <thrust/fill.h></span>
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-directive">#include <thrust/replace.h></span>
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-directive">#include <thrust/functional.h></span>
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-directive">#include <iostream></span>
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span> main(<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">void</span>)
{
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// allocate three device_vectors with 10 elements</span>
thrust::device_vector<<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span>> X(10);
thrust::device_vector<<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span>> Y(10);
thrust::device_vector<<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span>> Z(10);
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// initialize X to 0,1,2,3, ....</span>
thrust::sequence(X.begin(), X.end());
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// compute Y = -X</span>
thrust::transform(X.begin(), X.end(), Y.begin(), thrust::negate<<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span>>());
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// fill Z with twos</span>
thrust::fill(Z.begin(), Z.end(), 2);
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// compute Y = X mod 2</span>
thrust::transform(X.begin(), X.end(), Z.begin(), Y.begin(), thrust::modulus<<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span>>());
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// replace all the ones in Y with tens</span>
thrust::replace(Y.begin(), Y.end(), 1, 10);
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// print Y</span>
thrust::copy(Y.begin(), Y.end(), std::ostream_iterator<<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span>>(std::cout, <span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-string">"\n"</span>));
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">return</span> 0;
}
</pre><p class="p">While the functors in <samp class="ph codeph">thrust/functional.h</samp> cover most of the built-in arithmetic and comparison operations, we often want to do something different. For example, consider
the vector operation <samp class="ph codeph">y <- a * x + y</samp> where <samp class="ph codeph">x</samp> and <samp class="ph codeph">y</samp> are vectors and <samp class="ph codeph">a</samp> is a scalar constant. This is the well-known SAXPY operation provided by any <a class="xref" href="http://en.wikipedia.org/wiki/Basic_Linear_Algebra_Subprograms" target="_blank" shape="rect">BLAS</a> library.
</p>
<p class="p">If we want to implement SAXPY with Thrust we have a few options. The first is to use two transformations (one addition and
one multiplication) and a temporary vector filled with the value <samp class="ph codeph">a</samp>. A better choice is to use a single transformation with a user-defined functor that does exactly what we want. We illustrate
both approaches in the source code below.
</p><pre xml:space="preserve">
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">struct</span> saxpy_functor
{
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">const</span> <span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">float</span> a;
saxpy_functor(<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">float</span> _a) : a(_a) {}
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-attribute">__host__</span> <span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-attribute">__device__</span>
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">float</span> operator()(<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">const</span> <span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">float</span>& x, <span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">const</span> <span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">float</span>& y) <span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">const</span> {
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">return</span> a * x + y;
}
};
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">void</span> saxpy_fast(<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">float</span> A, thrust::device_vector<<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">float</span>>& X, thrust::device_vector<<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">float</span>>& Y)
{
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// Y <- A * X + Y</span>
thrust::transform(X.begin(), X.end(), Y.begin(), Y.begin(), saxpy_functor(A));
}
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">void</span> saxpy_slow(<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">float</span> A, thrust::device_vector<<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">float</span>>& X, thrust::device_vector<<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">float</span>>& Y)
{
thrust::device_vector<<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">float</span>> temp(X.size());
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// temp <- A</span>
thrust::fill(temp.begin(), temp.end(), A);
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// temp <- A * X</span>
thrust::transform(X.begin(), X.end(), temp.begin(), temp.begin(), thrust::multiplies<<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">float</span>>());
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// Y <- A * X + Y</span>
thrust::transform(temp.begin(), temp.end(), Y.begin(), Y.begin(), thrust::plus<<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">float</span>>());
}
</pre><p class="p">Both <samp class="ph codeph">saxpy_fast</samp> and <samp class="ph codeph">saxpy_slow</samp> are valid SAXPY implementations, however <samp class="ph codeph">saxpy_fast</samp> will be significantly faster than <samp class="ph codeph">saxpy_slow</samp>. Ignoring the cost of allocating the <samp class="ph codeph">temp</samp> vector and the arithmetic operations we have the following costs:
</p>
<p class="p"><samp class="ph codeph">fast_saxpy</samp>: performs 2N reads and N writes
</p>
<p class="p"><samp class="ph codeph">slow_saxpy</samp>: performs 4N reads and 3N writes
</p>
<p class="p">Since SAXPY is <dfn class="term">memory bound</dfn> (its performance is limited by memory bandwidth, not floating point performance) the larger number of reads and writes makes
<samp class="ph codeph">saxpy_slow</samp> much more expensive. In contrast, <samp class="ph codeph">saxpy_fast</samp> will perform about as fast as SAXPY in an optimized BLAS implementation. In memory bound algorithms like SAXPY it is generally
worthwhile to apply <dfn class="term">kernel fusion</dfn> (combining multiple operations into a single kernel) to minimize the number of memory transactions.
</p>
<p class="p"><samp class="ph codeph">thrust::transform</samp> only supports transformations with one or two input arguments (e.g.
<math xmlns="http://www.w3.org/1998/Math/MathML">
<mi>f</mi>
<mo stretchy="false">(</mo>
<mi>x</mi>
<mo stretchy="false">)</mo>
<mo stretchy="false">→</mo>
<mi>y</mi>
</math>
and
<math xmlns="http://www.w3.org/1998/Math/MathML">
<mi>f</mi>
<mo stretchy="false">(</mo>
<mi>x</mi>
<mo>,</mo>
<mi>x</mi>
<mo stretchy="false">)</mo>
<mo stretchy="false">→</mo>
<mi>y</mi>
</math>
). When a transformation uses more than two input arguments it is necessary to use a different approach. The <a class="xref" href="https://github.com/thrust/thrust/blob/master/examples/arbitrary_transformation.cu" target="_blank" shape="rect">arbitrary_transformation</a> example demonstrates a solution that uses <samp class="ph codeph">thrust::zip_iterator</samp> and <samp class="ph codeph">thrust::for_each</samp>.
</p>
</div>
</div>
<div class="topic concept nested1" xml:lang="en-us" id="reductions"><a name="reductions" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#reductions" name="reductions" shape="rect">3.2. Reductions</a></h3>
<div class="body conbody">
<p class="p">A reduction algorithm uses a binary operation to <dfn class="term">reduce</dfn> an input sequence to a single value. For example, the sum of an array of numbers is obtained by reducing the array with a
plus operation. Similarly, the maximum of an array is obtained by reducing with a operator that takes two inputs and returns
the maximum. The sum of an array is implemented with <samp class="ph codeph">thrust::reduce</samp> as follows:
</p><pre xml:space="preserve">
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span> sum = thrust::reduce(D.begin(), D.end(), (<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span>) 0, thrust::plus<<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span>>());
</pre><p class="p">The first two arguments to <samp class="ph codeph">reduce</samp> define the range of values while the third and fourth parameters provide the initial value and reduction operator respectively.
Actually, this kind of reduction is so common that it is the default choice when no initial value or operator is provided.
The following three lines are therefore equivalent:
</p><pre xml:space="preserve">
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span> sum = thrust::reduce(D.begin(), D.end(), (<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span>) 0, thrust::plus<<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span>>());
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span> sum = thrust::reduce(D.begin(), D.end(), (<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span>) 0);
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span> sum = thrust::reduce(D.begin(), D.end())
</pre><p class="p">Although <samp class="ph codeph">thrust::reduce</samp> is sufficient to implement a wide variety of reductions, Thrust provides a few additional functions for convenience (like
the STL). For example, <samp class="ph codeph">thrust::count</samp> returns the number of instances of a specific value in a given sequence:
</p><pre xml:space="preserve">
#include <thrust/count.h>
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-directive">#include <thrust/device_vector.h></span>
...
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// put three 1s in a device_vector</span>
thrust::device_vector<<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span>> vec(5,0);
vec[1] = 1;
vec[3] = 1;
vec[4] = 1;
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// count the 1s</span>
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span> result = thrust::count(vec.begin(), vec.end(), 1);
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// result is three</span>
</pre><p class="p">Other reduction operations include <samp class="ph codeph">thrust::count_if</samp>, <samp class="ph codeph">thrust::min_element</samp>, <samp class="ph codeph">thrust::max_element</samp>, <samp class="ph codeph">thrust::is_sorted</samp>, <samp class="ph codeph">thrust::inner_product</samp>, and several others. Refer to the <a class="xref" href="http://thrust.github.com/doc/group__reductions.html" target="_blank" shape="rect">documentation</a> for a complete listing.
</p>
<p class="p">The SAXPY example in the Transformations section showed how <dfn class="term">kernel fusion</dfn> can be used to reduce the number of memory transfers used by a transformation kernel. With <samp class="ph codeph">thrust::transform_reduce</samp> we can also apply kernel fusion to reduction kernels. Consider the following example which computes the <a class="xref" href="http://en.wikipedia.org/wiki/Norm_(mathematics)#Euclidean_norm" target="_blank" shape="rect">norm</a> of a vector.
</p><pre xml:space="preserve">
#include <thrust/transform_reduce.h>
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-directive">#include <thrust/functional.h></span>
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-directive">#include <thrust/device_vector.h></span>
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-directive">#include <thrust/host_vector.h></span>
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-directive">#include <cmath></span>
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// square<T> computes the square of a number f(x) -> x*x</span>
template <typename T>
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">struct</span> square
{
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-attribute">__host__</span> <span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-attribute">__device__</span>
T operator()(<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">const</span> T& x) <span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">const</span> {
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">return</span> x * x;
}
};
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span> main(<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">void</span>)
{
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// initialize host array</span>
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">float</span> x[4] = {1.0, 2.0, 3.0, 4.0};
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// transfer to device</span>
thrust::device_vector<<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">float</span>> d_x(x, x + 4);
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// setup arguments</span>
square<<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">float</span>> unary_op;
thrust::plus<<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">float</span>> binary_op;
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">float</span> init = 0;
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// compute norm</span>
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">float</span> norm = std::sqrt( thrust::transform_reduce(d_x.begin(), d_x.end(), unary_op, init, binary_op) );
std::cout << norm << std::endl;
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">return</span> 0;
}
</pre><p class="p">Here we have a <dfn class="term">unary operator</dfn> called <samp class="ph codeph">square</samp> that squares each element of the input sequence. The sum of squares is then computed using a standard <samp class="ph codeph">plus</samp> reduction. Like the slower version of the SAXPY transformation, we could implement <samp class="ph codeph">norm</samp> with multiple passes: first a <samp class="ph codeph">transform</samp> using <samp class="ph codeph">square</samp> or perhaps just <samp class="ph codeph">multiplies</samp> and then a <samp class="ph codeph">plus</samp> reduction over a temporary array. However this would be unnecessarily wasteful and considerably slower. By fusing the square
operation with the reduction kernel we again have a highly optimized implementation which offers the same performance as hand-written
kernels.
</p>
</div>
</div>
<div class="topic concept nested1" xml:lang="en-us" id="prefix-sums"><a name="prefix-sums" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#prefix-sums" name="prefix-sums" shape="rect">3.3. Prefix-Sums</a></h3>
<div class="body conbody">
<p class="p">Parallel prefix-sums, or <dfn class="term">scan</dfn> operations, are important building blocks in many parallel algorithms such as stream compaction and radix sort. Consider
the following source code which illustrates an <dfn class="term">inclusive scan</dfn> operation using the default <samp class="ph codeph">plus</samp> operator:
</p><pre xml:space="preserve">
#include <thrust/scan.h>
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span> data[6] = {1, 0, 2, 2, 1, 3};
thrust::inclusive_scan(data, data + 6, data); <span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// in-place scan</span>
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// data is now {1, 1, 3, 5, 6, 9}</span>
</pre><p class="p">In an inclusive scan each element of the output is the corresponding <a class="xref" href="http://en.wikipedia.org/wiki/Partial_sum#Basic_properties" target="_blank" shape="rect">partial sum</a> of the input range. For example, <samp class="ph codeph">data[2] = data[0] + data[1] + data[2]</samp>. An <dfn class="term">exclusive</dfn> scan is similar, but shifted by one place to the right:
</p><pre xml:space="preserve">
#include <thrust/scan.h>
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span> data[6] = {1, 0, 2, 2, 1, 3};
thrust::exclusive_scan(data, data + 6, data); <span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// in-place scan</span>
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// data is now {0, 1, 1, 3, 5, 6}</span>
</pre><p class="p">So now <samp class="ph codeph">data[2] = data[0] + data[1]</samp>. As these examples show, <samp class="ph codeph">inclusive_scan</samp> and <samp class="ph codeph">exclusive_scan</samp> are permitted to be performed in-place. Thrust also provides the functions <samp class="ph codeph">transform_inclusive_scan</samp> and <samp class="ph codeph">transform_exclusive_scan</samp> which apply a unary function to the input sequence before performing the scan. Refer to the <a class="xref" href="http://thrust.github.com/doc/group__prefixsums.html" target="_blank" shape="rect">documentation</a> for a complete list of scan variants.
</p>
</div>
</div>
<div class="topic concept nested1" xml:lang="en-us" id="reordering"><a name="reordering" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#reordering" name="reordering" shape="rect">3.4. Reordering</a></h3>
<div class="body conbody">
<p class="p">Thrust provides support for <dfn class="term">partitioning</dfn> and <dfn class="term">stream compaction</dfn> through the following algorithms:
</p>
<p class="p"><samp class="ph codeph">copy_if</samp> : copy elements that pass a predicate test
</p>
<p class="p"><samp class="ph codeph">partition</samp> : reorder elements according to a predicate (true values precede false values)
</p>
<p class="p"><samp class="ph codeph">remove</samp> and <samp class="ph codeph">remove_if</samp> : remove elements that fail a predicate test
</p>
<p class="p"><samp class="ph codeph">unique</samp>: remove consecutive duplicates within a sequence Refer to the <a class="xref" href="http://thrust.github.com/doc/group__reordering.html" target="_blank" shape="rect">documentation</a> for a complete list of reordering functions and examples of their usage.
</p>
</div>
</div>
<div class="topic concept nested1" xml:lang="en-us" id="sorting"><a name="sorting" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#sorting" name="sorting" shape="rect">3.5. Sorting</a></h3>
<div class="body conbody">
<p class="p">Thrust offers several functions to sort data or rearrange data according to a given criterion. The <samp class="ph codeph">thrust::sort</samp> and <samp class="ph codeph">thrust::stable_sort</samp> functions are direct analogs of <a class="xref" href="http://www.sgi.com/tech/stl/sort.html" target="_blank" shape="rect">sort</a> and <a class="xref" href="http://www.sgi.com/tech/stl/stable_sort.html" target="_blank" shape="rect">stable_sort</a> in the STL.
</p><pre xml:space="preserve">
#include <thrust/sort.h>
...
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">const</span> <span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span> N = 6;
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span> A[N] = {1, 4, 2, 8, 5, 7};
thrust::sort(A, A + N);
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// A is now {1, 2, 4, 5, 7, 8}</span>
</pre><p class="p">In addition, Thrust provides <samp class="ph codeph">thrust::sort_by_key</samp> and <samp class="ph codeph">thrust::stable_sort_by_key</samp>, which sort key-value pairs stored in separate places.
</p><pre xml:space="preserve">
#include <thrust/sort.h>
...
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">const</span> <span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span> N = 6;
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span> keys[N] = { 1, 4, 2, 8, 5, 7};
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">char</span> values[N] = {<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-string">'a'</span>, <span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-string">'b'</span>, <span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-string">'c'</span>, <span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-string">'d'</span>, <span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-string">'e'</span>, <span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-string">'f'</span>};
thrust::sort_by_key(keys, keys + N, values);
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// keys is now { 1, 2, 4, 5, 7, 8}</span>
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// values is now {'a', 'c', 'b', 'e', 'f', 'd'}</span>
</pre><p class="p">Like their STL brethren, the sorting functions also accept user-defined comparison operators:</p><pre xml:space="preserve">
#include <thrust/sort.h>
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-directive">#include <thrust/functional.h></span>
...
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">const</span> <span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span> N = 6;
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span> A[N] = {1, 4, 2, 8, 5, 7};
thrust::stable_sort(A, A + N, thrust::greater<<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span>>());
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// A is now {8, 7, 5, 4, 2, 1}</span>
</pre></div>
</div>
</div>
<div class="topic concept nested0" xml:lang="en-us" id="fancy-iterators"><a name="fancy-iterators" shape="rect">
<!-- --></a><h2 class="title topictitle1"><a href="#fancy-iterators" name="fancy-iterators" shape="rect">4. Fancy Iterators</a></h2>
<div class="body conbody">
<p class="p">Fancy iterators perform a variety of valuable purposes. In this section we’ll show how fancy iterators allow us to attack
a broader class problems with the standard Thrust algorithms. For those familiar with the <a class="xref" href="http://www.boost.org/" target="_blank" shape="rect">Boost C++ Library</a>, note that our fancy iterators were inspired by (and generally derived from) those in <a class="xref" href="http://www.boost.org/doc/libs/1_40_0/libs/iterator/doc/index.html" target="_blank" shape="rect">Boost Iterator Library</a>.
</p>
</div>
<div class="topic concept nested1" xml:lang="en-us" id="constant-iterator"><a name="constant-iterator" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#constant-iterator" name="constant-iterator" shape="rect">4.1. constant_iterator</a></h3>
<div class="body conbody">
<p class="p">Arguably the simplest of the bunch, <samp class="ph codeph">constant_iterator</samp> is simply an iterator that returns the same value whenever we access it. In the following example we initialize a constant
iterator with the value 10.
</p><pre xml:space="preserve">
#include <thrust/iterator/constant_iterator.h>
...
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// create iterators</span>
thrust::constant_iterator<<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span>> first(10);
thrust::constant_iterator<<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span>> last = first + 3;
first[0] <span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// returns 10</span>
first[1] <span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// returns 10</span>
first[100] <span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// returns 10</span>
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// sum of [first, last)</span>
thrust::reduce(first, last); <span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// returns 30 (i.e. 3 * 10)</span>
</pre><p class="p">Whenever an input sequence of constant values is needed, <samp class="ph codeph">constant_iterator</samp> is a convenient and efficient solution.
</p>
</div>
</div>
<div class="topic concept nested1" xml:lang="en-us" id="counting-iterator"><a name="counting-iterator" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#counting-iterator" name="counting-iterator" shape="rect">4.2. counting_iterator</a></h3>
<div class="body conbody">
<p class="p">If a sequence of increasing values is required, then <samp class="ph codeph">counting_iterator</samp> is the appropriate choice. Here we initialize a <samp class="ph codeph">counting_iterator</samp> with the value 10 and access it like an array.
</p><pre xml:space="preserve">
#include <thrust/iterator/counting_iterator.h>
...
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// create iterators</span>
thrust::counting_iterator<<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span>> first(10);
thrust::counting_iterator<<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span>> last = first + 3;
first[0] <span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// returns 10</span>
first[1] <span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// returns 11</span>
first[100] <span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// returns 110</span>
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// sum of [first, last)</span>
thrust::reduce(first, last); <span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// returns 33 (i.e. 10 + 11 + 12)</span>
</pre><p class="p">While <samp class="ph codeph">constant_iterator</samp> and <samp class="ph codeph">counting_iterator</samp> act as arrays, they don’t actually require any memory storage. Whenever we dereference one of these iterators it generates
the appropriate value on-the-fly and returns it to the calling function.
</p>
</div>
</div>
<div class="topic concept nested1" xml:lang="en-us" id="transform-iterator"><a name="transform-iterator" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#transform-iterator" name="transform-iterator" shape="rect">4.3. transform_iterator</a></h3>
<div class="body conbody">
<p class="p">In the Algorithms section we spoke about kernel fusion, i.e. combining separate algorithms like <samp class="ph codeph">transform</samp> and <samp class="ph codeph">reduce</samp> into a single <samp class="ph codeph">transform_reduce</samp> operation. The <samp class="ph codeph">transform_iterator</samp> allows us to apply the same technique, even when we don’t have a special <samp class="ph codeph">transform_xxx</samp> version of the algorithm. This example shows another way to fuse a transformation with a reduction, this time with just plain
<samp class="ph codeph">reduce</samp> applied to a <samp class="ph codeph">transform_iterator</samp>.
</p><pre xml:space="preserve">
#include <thrust/iterator/transform_iterator.h>
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// initialize vector</span>
thrust::device_vector<<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span>> vec(3);
vec[0] = 10; vec[1] = 20; vec[2] = 30;
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// create iterator (type omitted)</span>
...
first = thrust::make_transform_iterator(vec.begin(), negate<<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span>>());
...
last = thrust::make_transform_iterator(vec.end(), negate<<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span>>());
first[0] <span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// returns -10</span>
first[1] <span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// returns -20</span>
first[2] <span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// returns -30</span>
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// sum of [first, last)</span>
thrust::reduce(first, last); <span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// returns -60 (i.e. -10 + -20 + -30)</span>
</pre><p class="p">Note, we have omitted the types for iterators <samp class="ph codeph">first</samp> and <samp class="ph codeph">last</samp> for simplicity. One downside of <samp class="ph codeph">transform_iterator</samp> is that it can be cumbersome to specify the full type of the iterator, which can be quite lengthy. For this reason, it is
common practice to simply put the call to <samp class="ph codeph">make_transform_iterator</samp> in the arguments of the algorithm being invoked. For example,
</p><pre xml:space="preserve">
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// sum of [first, last)</span>
thrust::reduce(thrust::make_transform_iterator(vec.begin(), negate<<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span>>()),
thrust::make_transform_iterator(vec.end(), negate<<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span>>()));
</pre><p class="p">allows us to avoid creating a variable to store <samp class="ph codeph">first</samp> and <samp class="ph codeph">last</samp>.
</p>
</div>
</div>
<div class="topic concept nested1" xml:lang="en-us" id="permutation-iterator"><a name="permutation-iterator" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#permutation-iterator" name="permutation-iterator" shape="rect">4.4. permutation_iterator</a></h3>
<div class="body conbody">
<p class="p">In the previous section we showed how <samp class="ph codeph">transform_iterator</samp> is used to fuse a transformation with another algorithm to avoid unnecessary memory operations. The <samp class="ph codeph">permutation_iterator</samp> is similar: it allows us to fuse gather and scatter operations with Thrust algorithms, or even other fancy iterators. The
following example shows how to fuse a gather operation with a reduction.
</p><pre xml:space="preserve">
#include <thrust/iterator/permutation_iterator.h>
...
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// gather locations</span>
thrust::device_vector<<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span>> map(4);
map[0] = 3;
map[1] = 1;
map[2] = 0;
map[3] = 5;
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// array to gather from</span>
thrust::device_vector<<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span>> source(6);
source[0] = 10;
source[1] = 20;
source[2] = 30;
source[3] = 40;
source[4] = 50;
source[5] = 60;
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// fuse gather with reduction: </span>
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// sum = source[map[0]] + source[map[1]] + ...</span>
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span> sum = thrust::reduce(thrust::make_permutation_iterator(source.begin(), map.begin()),
thrust::make_permutation_iterator(source.begin(), map.end()));
</pre><p class="p">Here we have used the <samp class="ph codeph">make_permutation_iterator</samp> function to simplify the construction of the <samp class="ph codeph">permutation_iterator</samp>s. The first argument to <samp class="ph codeph">make_permutation_iterator</samp> is the source array of the gather operation and the second is the list of map indices. Note that we pass in <samp class="ph codeph">source.begin()</samp> for the first argument in both cases, but vary the second argument to define the beginning and end of the sequence.
</p>
<p class="p">When a <samp class="ph codeph">permutation_iterator</samp> is used as an output sequence of a function it is equivalent to fusing a scatter operation to the algorithm. In general <samp class="ph codeph">permutation_iterator</samp> allows you to operate on specific set of values in a sequence instead of the entire sequence.
</p>
</div>
</div>
<div class="topic concept nested1" xml:lang="en-us" id="zip-iterator"><a name="zip-iterator" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#zip-iterator" name="zip-iterator" shape="rect">4.5. zip_iterator</a></h3>
<div class="body conbody">
<p class="p">Keep reading, we’ve saved the best iterator for last! The <samp class="ph codeph">zip_iterator</samp> is an extremely useful gadget: it takes multiple input sequences and yields a sequence of tuples. In this example we “zip”
together a sequence of <samp class="ph codeph">int</samp> and a sequence of <samp class="ph codeph">char</samp> into a sequence of <samp class="ph codeph">tuple<int,char></samp> and compute the <samp class="ph codeph">tuple</samp> with the maximum value.
</p><pre xml:space="preserve">
#include <thrust/iterator/zip_iterator.h>
...
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// initialize vectors</span>
thrust::device_vector<<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span>> A(3);
thrust::device_vector<<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">char</span>> B(3);
A[0] = 10; A[1] = 20; A[2] = 30;
B[0] = <span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-string">'x'</span>; B[1] = <span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-string">'y'</span>; B[2] = <span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-string">'z'</span>;
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// create iterator (type omitted)</span>
first = thrust::make_zip_iterator(thrust::make_tuple(A.begin(), B.begin()));
last = thrust::make_zip_iterator(thrust::make_tuple(A.end(), B.end()));
first[0] <span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// returns tuple(10, 'x')</span>
first[1] <span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// returns tuple(20, 'y')</span>
first[2] <span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// returns tuple(30, 'z')</span>
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// maximum of [first, last)</span>
thrust::maximum< tuple<<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span>,<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">char</span>> > binary_op;
thrust::tuple<<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span>,<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">char</span>> init = first[0];
thrust::reduce(first, last, init, binary_op); <span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-comment">// returns tuple(30, 'z')</span>
</pre><p class="p">What makes <samp class="ph codeph">zip_iterator</samp> so useful is that most algorithms accept either one, or occasionally two, input sequences. The <samp class="ph codeph">zip_iterator</samp> allows us to combine many independent sequences into a single sequence of tuples, which can be processed by a broad set of
algorithms.
</p>
<p class="p">Refer to the <a class="xref" href="https://github.com/thrust/thrust/blob/master/examples/arbitrary_transformation.cu" target="_blank" shape="rect">arbitrary_transformation</a> example to see how to implement a ternary transformation with <samp class="ph codeph">zip_iterator</samp> and <samp class="ph codeph">for_each</samp>. A simple extension of this example would allow you to compute transformations with multiple <dfn class="term">output</dfn> sequences as well.
</p>
<p class="p">In addition to convenience, <samp class="ph codeph">zip_iterator</samp> allows us to implement programs more efficiently. For example, storing 3d points as an array of <samp class="ph codeph">float3</samp> in CUDA is generally a bad idea, since array accesses are not properly coalesced. With <samp class="ph codeph">zip_iterator</samp> we can store the three coordinates in three separate arrays, which does permit coalesced memory access. In this case, we
use <samp class="ph codeph">zip_iterator</samp> to create a <dfn class="term">virtual</dfn> array of 3d vectors which we can feed in to Thrust algorithms. Refer to the <a class="xref" href="https://github.com/thrust/thrust/blob/master/examples/dot_products_with_zip.cu" target="_blank" shape="rect">dot_products_with_zip</a> example for additional details.
</p>
</div>
</div>
</div>
<div class="topic concept nested0" xml:lang="en-us" id="additional-resources"><a name="additional-resources" shape="rect">
<!-- --></a><h2 class="title topictitle1"><a href="#additional-resources" name="additional-resources" shape="rect">5. Additional Resources</a></h2>
<div class="body conbody">
<p class="p">This guide only scratches the surface of what you can do with Thrust. The following resources can help you learn to do more
with Thrust or provide assistance when problems arise.
</p>
<p class="p">Comprehensive <a class="xref" href="https://github.com/thrust/thrust/wiki/Documentation" target="_blank" shape="rect">documentation</a> of Thrust’s API
</p>
<p class="p">A list of <a class="xref" href="https://github.com/thrust/thrust/wiki/Frequently-Asked-Questions" target="_blank" shape="rect">Frequently Asked Questions</a></p>
<p class="p">Collection of <a class="xref" href="https://github.com/thrust/thrust/tree/master/examples" target="_blank" shape="rect">example</a> programs
</p>
<p class="p">We strongly encourage users to subscribe to the <a class="xref" href="http://groups.google.com/group/thrust-users" target="_blank" shape="rect">thrust-users</a> mailing list. The mailing list is a great place to seek out help from the Thrust developers and other Thrust users.
</p>
</div>
</div>
<div class="topic concept nested0" id="notices-header"><a name="notices-header" shape="rect">
<!-- --></a><h2 class="title topictitle1"><a href="#notices-header" name="notices-header" shape="rect">Notices</a></h2>
<div class="topic reference nested1" id="notice"><a name="notice" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#notice" name="notice" shape="rect"></a></h3>
<div class="body refbody">
<div class="section">
<h3 class="title sectiontitle">Notice</h3>
<p class="p">ALL NVIDIA DESIGN SPECIFICATIONS, REFERENCE BOARDS, FILES, DRAWINGS, DIAGNOSTICS, LISTS, AND OTHER DOCUMENTS (TOGETHER AND
SEPARATELY, "MATERIALS") ARE BEING PROVIDED "AS IS." NVIDIA MAKES NO WARRANTIES, EXPRESSED, IMPLIED, STATUTORY, OR OTHERWISE
WITH RESPECT TO THE MATERIALS, AND EXPRESSLY DISCLAIMS ALL IMPLIED WARRANTIES OF NONINFRINGEMENT, MERCHANTABILITY, AND FITNESS
FOR A PARTICULAR PURPOSE.
</p>
<p class="p">Information furnished is believed to be accurate and reliable. However, NVIDIA Corporation assumes no responsibility for the
consequences of use of such information or for any infringement of patents or other rights of third parties that may result
from its use. No license is granted by implication of otherwise under any patent rights of NVIDIA Corporation. Specifications
mentioned in this publication are subject to change without notice. This publication supersedes and replaces all other information
previously supplied. NVIDIA Corporation products are not authorized as critical components in life support devices or systems
without express written approval of NVIDIA Corporation.
</p>
</div>
</div>
</div>
<div class="topic reference nested1" id="trademarks"><a name="trademarks" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#trademarks" name="trademarks" shape="rect"></a></h3>
<div class="body refbody">
<div class="section">
<h3 class="title sectiontitle">Trademarks</h3>
<p class="p">NVIDIA and the NVIDIA logo are trademarks or registered trademarks of NVIDIA Corporation
in the U.S. and other countries. Other company and product names may be trademarks of
the respective companies with which they are associated.
</p>
</div>
</div>
</div>
<div class="topic reference nested1" id="copyright-past-to-present"><a name="copyright-past-to-present" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#copyright-past-to-present" name="copyright-past-to-present" shape="rect"></a></h3>
<div class="body refbody">
<div class="section">
<h3 class="title sectiontitle">Copyright</h3>
<p class="p">© <span class="ph">2011</span>-<span class="ph">2014</span> NVIDIA
Corporation. All rights reserved.
</p>
<p class="p">This product includes software developed by the Syncro Soft SRL (http://www.sync.ro/).</p>
</div>
</div>
</div>
</div>
<hr id="contents-end"></hr>
</article>
</div>
</div>
<script language="JavaScript" type="text/javascript" charset="utf-8" src="../common/formatting/common.min.js"></script>
<script language="JavaScript" type="text/javascript" charset="utf-8" src="../common/scripts/google-analytics/google-analytics-write.js"></script>
<script language="JavaScript" type="text/javascript" charset="utf-8" src="../common/scripts/google-analytics/google-analytics-tracker.js"></script>
<script type="text/javascript">var switchTo5x=true;</script><script type="text/javascript" src="http://w.sharethis.com/button/buttons.js"></script><script type="text/javascript">stLight.options({publisher: "998dc202-a267-4d8e-bce9-14debadb8d92", doNotHash: false, doNotCopy: false, hashAddressBar: false});</script></body>
</html>
distrib
>
Mageia
>
5
>
x86_64
>
media
>
nonfree-updates
>
by-pkgid
>
fd8445e7e4d58b8cfe6e0150bd441ee1
>
files
>
1736
