thrust快速入门指南（并行算法库，类似C++的STL）

thrust快速入门指南

Introduction

  Thrust是基于标准模板库（STL）的并行平台的C ++模板库。Thrust允许您通过高级接口以最少的编程工作实现高性能并行应用程序，该接口可与C ++，CUDA，OpenMP和TBB等技术完全互操作。
  Thrust提供了丰富的数据并行原语集合，例如扫描，排序和缩减，它们可以组合在一起，通过简洁易读的源代码实现复杂的算法。通过根据这些高级抽象描述您的计算，您可以为Thrust提供自动选择最有效实现的自由。因此，Thrust可用于CUDA应用程序的快速原型设计，其中程序员生产力最重要，而且在生产中，稳健性和绝对性能至关重要。
  本文档描述了如何使用Thrust开发并行应用程序。即使您具有有限的C ++或并行编程经验，也可以访问本教程。

Prerequisites

  Thrust v1.6.0与CUDA 4.1（首选）和CUDA 4.0兼容。您可以通过nvcc --version在命令行上运行来确认已安装CUDA 。例如，在Linux系统上，

$ nvcc --version
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2011 NVIDIA Corporation
Built on Thu_Jan_12_14:41:45_PST_2012
Cuda compilation tools, release 4.1, V0.2.1221

  如果您使用的是CUDA 4.0或更高版本，那么您的系统上已经安装了Thrust，您可以安全地跳到下一部分。
  由于Thrust是一个C ++模板库，因此无需“构建”。只需从下载部分下载最新版本，然后将zip文件的内容解压缩到一个目录中。我们建议将Thrust安装到CUDA include目录中，这通常是
  /usr/local/cuda/include/ 在Linux和Mac OSX上
  C:\CUDA\include\ 在Windows系统上
如果您无法将Thrust安装到CUDA include目录，那么您可以将Thrust放在主目录中的某个位置，例如：/home/nathan/libraries/。
  警告：不要将Thrust安装到标准包含路径之类的/usr/local/include/。似乎nvcc对这些路径的处理方式不同于上面的建议，这会导致错误error: expected primary-expression before ‘<’ token.

Simple Example

  让我们用Thrust编译一个简单的程序，以确保满足所有先决条件。将以下源代码保存到名为的文件中version.cu。

 #include <thrust/version.h>
 #include <iostream>
int main(void)
{
  int major = THRUST_MAJOR_VERSION;
  int minor = THRUST_MINOR_VERSION;

  std::cout << "Thrust v" << major << "." << minor << std::endl;

  return 0;
}

  现在编译version.cu带nvcc。如果将Thrust安装到CUDA include目录，那么以下命令应该可以正常工作。

$ ls
thrust  version.cu
$ nvcc version.cu -o version
$ ls
thrust  version version.cu
$ ./version
Thrust v1.6

  如果Thrust 目录放在其他地方，请使用该-I选项告知nvcc要查看的位置。例如，如果放入推力，/home/nathan/libraries/则应使用以下命令。

$ nvcc version.cu -o version -I /home/nathan/libraries/

Vectors

  Thrust提供两个矢量容器，host_vector和device_vector。顾名思义，host_vector存储在CPU的系统或“主机”内存中，同时device_vector存在于GPU的“设备”内存中。Thrust的向量容器就像std::vector在C ++标准库中一样。喜欢std::vector，host_vector并且device_vector是可以动态调整大小的通用容器（能够存储任何数据类型）。以下源代码说明了Thrust的向量容器的使用。

 #include <thrust/host_vector.h>
 #include <thrust/device_vector.h>

 #include <iostream>

int main(void)
{
  // H has storage for 4 integers
  thrust::host_vector<int> H(4);

  // initialize individual elements
  H[0] = 14;
  H[1] = 20;
  H[2] = 38;
  H[3] = 46;
    
  // H.size() returns the size of vector H
  std::cout << "H has size " << H.size() << std::endl;

  // print contents of H
  for(int i = 0; i < H.size(); i++)
  {
    std::cout << "H[" << i << "] = " << H[i] << std::endl;
  }

  // resize H
  H.resize(2);
    
  std::cout << "H now has size " << H.size() << std::endl;

  // Copy host_vector H to device_vector D
  thrust::device_vector<int> D = H;
    
  // elements of D can be modified
  D[0] = 99;
  D[1] = 88;
    
  // print contents of D
  for(int i = 0; i < D.size(); i++)
  {
    std::cout << "D[" << i << "] = " << D[i] << std::endl;
  }

  // H and D are automatically destroyed when the function returns
  return 0;
}

  如此示例所示，=操作员可用于将a复制host_vector到a device_vector（反之亦然）。该=运营商还可以用来复制host_vector到host_vector或device_vector到device_vector。另请注意，device_vector可以使用标准括号表示法访问a的各个元素。但是，因为这些访问中的每一个都需要调用cudaMemcpy，所以应该谨慎使用它们。稍后我们将介绍一些更有效的技术。
  将向量的所有元素初始化为特定值或仅将一组值从一个向量复制到另一个向量通常很有用。Thrust提供了一些方法来执行这些操作。

#include <thrust/host_vector.h>
#include <thrust/device_vector.h>

#include <thrust/copy.h>
#include <thrust/fill.h>
#include <thrust/sequence.h>

#include <iostream>

int main(void)
{
  // initialize all ten integers of a device_vector to 1
  thrust::device_vector<int> D(10, 1);

  // set the first seven elements of a vector to 9
  thrust::fill(D.begin(), D.begin() + 7, 9);

  // initialize a host_vector with the first five elements of D
  thrust::host_vector<int> H(D.begin(), D.begin() + 5);

  // set the elements of H to 0, 1, 2, 3, ...
  thrust::sequence(H.begin(), H.end());

  // copy all of H back to the beginning of D
  thrust::copy(H.begin(), H.end(), D.begin());

  // print D
  for(int i = 0; i < D.size(); i++)
  {
    std::cout << "D[" << i << "] = " << D[i] << std::endl;
  }

  return 0;
}

  这里我们说明使用的fill，copy和sequence功能。该copy函数可用于将一系列主机或设备元素复制到另一个主机或设备向量。与相应的C ++标准库函数一样，thrust::fill只需将元素范围设置为特定值即可。Thrust的sequence函数可用于创建一系列等间距值。

Thrust Namespace

  你会发现，我们使用的东西，像thrust::host_vector或thrust::copy在我们的例子。该thrust::部分告诉C ++编译器我们想要查看thrust特定函数或类的命名空间。命名空间是避免名称冲突的好方法。例如，与C ++标准库中提供的thrust::copy不同std::copy。C ++命名空间允许我们区分这两个copy函数。

Iterators and Static Dispatching

  在本节中，我们使用像H.begin()和等H.end()偏移的表达式D.begin() + 7。结果begin()和end()在C ++中称为迭代器。对于矢量容器（实际上只是数组），迭代器可以被认为是指向数组元素的指针。因此，H.begin()是一个迭代器，它指向存储在H向量内的数组的第一个元素。类似地，H.end()指向元素一个超过H向量的最后一个元素。
  虽然向量迭代器与指针类似，但它们带有更多信息。请注意，我们不必告诉thrust::fill它在device_vector迭代器上运行。此信息以返回的迭代器类型捕获，该迭代器的类型与返回D.begin()的类型不同H.begin()。调用Thrust函数时，它会检查迭代器的类型，以确定是使用主机还是设备实现。此过程称为静态分派，因为主机/设备分派在编译时解析。请注意，这意味着调度进程没有运行时开销。
  您可能想知道当“原始”指针用作Thrust函数的参数时会发生什么。与STL一样，Thrust允许这种用法，它将调度算法的主机路径。如果有问题的指针实际上是指向设备内存的指针，那么thrust::device_ptr在调用函数之前，您需要将其包装起来。例如：

size_t N = 10;

// raw pointer to device memory
int * raw_ptr;
cudaMalloc((void **) &raw_ptr, N * sizeof(int));

// wrap raw pointer with a device_ptr 
thrust::device_ptr<int> dev_ptr(raw_ptr);

// use device_ptr in thrust algorithms
thrust::fill(dev_ptr, dev_ptr + N, (int) 0);

  要从device_ptr中提取原始指针，raw_pointer_cast应按如下方式应用：

size_t N = 10;

// create a device_ptr 
thrust::device_ptr<int> dev_ptr = thrust::device_malloc<int>(N);
     
// extract raw pointer from device_ptr
int * raw_ptr = thrust::raw_pointer_cast(dev_ptr);

  区分迭代器和指针的另一个原因是迭代器可用于遍历多种数据结构。例如，C ++标准库提供了一个链表list（std::list），它提供了双向（但不是随机访问）迭代器。尽管Thrust不提供此类容器的设备实现，但它与它们兼容。

#include <thrust/device_vector.h>
#include <thrust/copy.h>
#include <list>
#include <vector>

int main(void)
{
  // create an STL list with 4 values
  std::list<int> stl_list;

  stl_list.push_back(10);
  stl_list.push_back(20);
  stl_list.push_back(30);
  stl_list.push_back(40);

  // initialize a device_vector with the list
  thrust::device_vector<int> D(stl_list.begin(), stl_list.end());

  // copy a device_vector into an STL vector
  std::vector<int> stl_vector(D.size());
  thrust::copy(D.begin(), D.end(), stl_vector.begin());

  return 0;
}

  对于未来参考：我们到目前为止所涉及的迭代器是有用的，但相当基础。除了这些普通的迭代器之外，Thrust还提供了一组名称为counting_iterator和的花式迭代器zip_iterator。虽然它们看起来和感觉像普通的迭代器，但是花哨的迭代器能够提供更多令人兴奋的东西。我们将在本教程后面重新讨论这个主题。

Algorithms

  Thrust提供了大量常见的并行算法。其中许多算法在C ++标准库中都有直接类比，当存在等效的标准库函数时，我们选择名称（例如thrust::sort和std::sort）。
  Thrust中的所有算法都具有主机和设备的实现。具体来说，当使用主机迭代器调用Thrust算法时，将调度主机路径。类似地，当使用设备迭代器定义范围时，将调用设备实现。
  除了thrust::copy可以在主机和设备之间复制数据之外，Thrust算法的所有迭代器参数都应该位于同一位置：要么全部在主机上，要么全部在设备上。违反此要求时，编译器将生成错误消息。

Transformations

  转换是将操作应用于一组（零个或多个）输入范围中的每个元素，然后将结果存储在目标范围内的算法。我们已经看到的一个例子是thrust::fill，它将范围的所有元素设置为指定值。其他转变包括thrust::sequence，thrust::replace当然thrust::transform。有关完整列表，请参阅文档。
  以下源代码演示了几种转换算法。请注意，thrust::negate与thrust::modulus已知为函子在C ++术语。推力提供这些和其他常见的仿函数像plus和multiplies文件中thrust/functional.h。

#include <thrust/device_vector.h>
#include <thrust/transform.h>
#include <thrust/sequence.h>
#include <thrust/copy.h>
#include <thrust/fill.h>
#include <thrust/replace.h>
#include <thrust/functional.h>
#include <iostream>

int main(void)
{
  // allocate three device_vectors with 10 elements
  thrust::device_vector<int> X(10);
  thrust::device_vector<int> Y(10);
  thrust::device_vector<int> Z(10);

  // initialize X to 0,1,2,3, ....
  thrust::sequence(X.begin(), X.end());

  // compute Y = -X
  thrust::transform(X.begin(), X.end(), Y.begin(), thrust::negate<int>());

  // fill Z with twos
  thrust::fill(Z.begin(), Z.end(), 2);

  // compute Y = X mod 2
  thrust::transform(X.begin(), X.end(), Z.begin(), Y.begin(), thrust::modulus<int>());

  // replace all the ones in Y with tens
  thrust::replace(Y.begin(), Y.end(), 1, 10);

  // print Y
  thrust::copy(Y.begin(), Y.end(), std::ostream_iterator<int>(std::cout, "\n"));
   
  return 0;    
}

  虽然thrust/functional.h算子涵盖了大多数内置算术和比较操作，但我们经常想做一些不同的事情。例如，考虑向量运算y <- a * x + ywhere x和y是向量并且a是标量常量。这是任何BLAS库提供的众所周知的SAXPY操作。
  如果我们想用Thrust实现SAXPY，我们有几个选择。第一种是使用两个转换（一个加法和一个乘法）和一个填充值a的临时向量。更好的选择是使用一个用户定义的仿函数进行单一转换，它完全符合我们的要求。我们在下面的源代码中说明了这两种方法。

struct saxpy_functor
{
  const float a;

  saxpy_functor(float _a) : a(_a) {}

  __host__ __device__
  float operator()(const float& x, const float& y) const
  { 
    return a * x + y;
  }
};

void saxpy_fast(float A, thrust::device_vector<float>& X, thrust::device_vector<float>& Y)
{
  // Y <- A * X + Y
  thrust::transform(X.begin(), X.end(), Y.begin(), Y.begin(), saxpy_functor(A));
}

void saxpy_slow(float A, thrust::device_vector<float>& X, thrust::device_vector<float>& Y)
{
  thrust::device_vector<float> temp(X.size());
   
  // temp <- A
  thrust::fill(temp.begin(), temp.end(), A);
    
  // temp <- A * X
  thrust::transform(X.begin(), X.end(), temp.begin(), temp.begin(), thrust::multiplies<float>());

  // Y <- A * X + Y
  thrust::transform(temp.begin(), temp.end(), Y.begin(), Y.begin(), thrust::plus<float>());
}

  这两个saxpy_fast和saxpy_slow有效SAXPY实现，但是saxpy_fast会比显著快saxpy_slow。忽略分配临时向量和算术运算的成本，我们有以下成本：
  fast_saxpy：执行2N读取和N次写入
  slow_saxpy：执行4N读取和3N写入
  由于SAXPY 受内存限制（其性能受内存带宽限制，而非浮点性能），因此读取和写入的saxpy_slow数量越多，成本就越高。相比之下，saxpy_fast在优化的BLAS实现中，其执行速度与SAXPY一样快。在像SAXPY这样的内存绑定算法中，通常值得应用内核融合（将多个操作组合到单个内核中）以最小化内存事务的数量。
  thrust::transform仅支持具有一个或两个输入参数的转换（例如f(x) -> y和f(x,y) -> z）。当转换使用两个以上的输入参数时，必须使用不同的方法。该arbitrary_transformation示例演示了使用thrust::zip_iterator和的解决方案thrust::for_each。

Reductions

  缩减算法使用二进制运算将输入序列减少为单个值。例如，通过使用加号运算减少数组来获得数字数组的总和。类似地，通过使用两个输入并返回最大值的运算符进行减少来获得数组的最大值。数组的总和thrust::reduce如下实现：

int sum = thrust::reduce(D.begin(), D.end(), (int) 0, thrust::plus<int>());

  reduce定义值范围的前两个参数，而第三个和第四个参数分别提供初始值和减少运算符。实际上，这种减少是如此常见，以至于在没有提供初始值或运算符时它是默认选择。因此以下三行是等效的：

int sum = thrust::reduce(D.begin(), D.end(), (int) 0, thrust::plus<int>());
int sum = thrust::reduce(D.begin(), D.end(), (int) 0);
int sum = thrust::reduce(D.begin(), D.end());

  虽然thrust::reduce足以实现各种各样的约简，但Thrust提供了一些额外的功能以方便使用（如C ++标准库）。例如，thrust::count返回给定序列中特定值的实例数：

#include <thrust/count.h>
#include <thrust/device_vector.h>
...
// put three 1s in a device_vector
thrust::device_vector<int> vec(5,0);
vec[1] = 1;
vec[3] = 1;
vec[4] = 1;

// count the 1s
int result = thrust::count(vec.begin(), vec.end(), 1);
// result is three

  其他降低操作包括thrust::count_if，thrust::min_element，thrust::max_element，thrust::is_sorted，thrust::inner_product，和其他几个人。有关完整列表，请参阅文档。
  转换部分中的SAXPY示例显示了如何使用内核融合来减少转换内核使用的内存传输次数。随着thrust::transform_reduce我们也可以内核融合适用于减少内核。考虑以下用于计算向量范数的示例。

#include <thrust/transform_reduce.h>
#include <thrust/functional.h>
#include <thrust/device_vector.h>
#include <thrust/host_vector.h>
#include <cmath>

// square<T> computes the square of a number f(x) -> x*x
template <typename T>
struct square
{
  __host__ __device__
  T operator()(const T& x) const
  { 
    return x * x;
  }
};

int main(void)
{
  // initialize host array
  float x[4] = {1.0, 2.0, 3.0, 4.0};

  // transfer to device
  thrust::device_vector<float> d_x(x, x + 4);

  // setup arguments
  square<float>        unary_op;
  thrust::plus<float> binary_op;
  float init = 0;

  // compute norm
  float norm = std::sqrt( thrust::transform_reduce(d_x.begin(), d_x.end(), unary_op, init, binary_op) );

  std::cout << norm << std::endl;

  return 0;
}

  这里我们有一个一元运算符，它调用square输入序列的每个元素。然后使用标准plus缩减来计算平方和。与SAXPY转换的较慢版本一样，我们可以实现norm多次传递：首先transform使用square或者只是multiplies然后plus减少临时数组。然而，这将是不必要的浪费并且相当慢。通过将square操作与还原内核融合，我们再次具有高度优化的实现，其提供与手写内核相同的性能。

Prefix-Sums

  并行前缀和或扫描操作是许多并行算法中的重要构建块，例如流压缩和基数排序。请考虑以下源代码，该代码说明了使用default plus运算符的包容性扫描操作：

#include <thrust/scan.h>

int data[6] = {1, 0, 2, 2, 1, 3};

thrust::inclusive_scan(data, data + 6, data); // in-place scan

// data is now {1, 1, 3, 5, 6, 9}

  在包含扫描中，输出的每个元素是输入范围的对应部分和。例如，data[2] = data[0] + data[1] + data[2]。一个独特的扫描是相似的，但一个地方的右移：

#include <thrust/scan.h>

int data[6] = {1, 0, 2, 2, 1, 3};

thrust::exclusive_scan(data, data + 6, data); // in-place scan

// data is now {0, 1, 1, 3, 5, 6}

  所以现在data[2] = data[0] + data[1]。如这些示例所示，inclusive_scan并且exclusive_scan允许就地执行。Thrust还提供功能，transform_inclusive_scan并transform_exclusive_scan在执行扫描之前将一元函数应用于输入序列。有关扫描变体的完整列表，请参阅文档。

Reordering

  Thrust 通过以下算法为分区和流压缩提供支持：
  copy_if ：复制传递谓词测试的元素
  partition：根据谓词重新排序元素（true值在false值之前）
  remove和remove_if：删除未通过谓词测试的元素
  unique：删除范围内的连续重复项
  有关重新排序功能的完整列表及其用法示例，请参阅文档。

Sorting

  Thrust提供了几种根据给定标准对数据进行排序或重新排列数据的功能。该thrust::sort和thrust::stable_sort功能的直接类似物sort和stable_sortC ++标准资源库中

#include <thrust/sort.h>
...
const int N = 6;
int A[N] = {1, 4, 2, 8, 5, 7};
thrust::sort(A, A + N);
// A is now {1, 2, 4, 5, 7, 8}

  此外，Thrust提供thrust::sort_by_key和thrust::stable_sort_by_key存储在不同位置的键值对。

#include <thrust/sort.h>
...
const int N = 6;
int    keys[N] = {  1,   4,   2,   8,   5,   7};
char values[N] = {'a', 'b', 'c', 'd', 'e', 'f'};
thrust::sort_by_key(keys, keys + N, values);
// keys is now   {  1,   2,   4,   5,   7,   8}
// values is now {'a', 'c', 'b', 'e', 'f', 'd'}

  与他们的标准库兄弟一样，排序函数也接受用户定义的比较运算符：

#include <thrust/sort.h>
#include <thrust/functional.h>
...
const int N = 6;
int A[N] = {1, 4, 2, 8, 5, 7};
thrust::stable_sort(A, A + N, thrust::greater<int>());
// A is now {8, 7, 5, 4, 2, 1}

Fancy Iterators

  花式迭代器执行各种有价值的目的。在本节中，我们将展示花哨的迭代器如何允许我们使用标准Thrust算法攻击更广泛的问题。对于那些熟悉Boost C ++库的人来说，请注意我们的花式迭代器的灵感来自Boost迭代器库中的（并且通常来自）。

constant_iterator

  可以说是最简单的一堆constant_iterator迭代器，只要我们取消引用它就会返回相同的值。在下面的示例中，我们constant_iterator使用值初始化a 10。

#include <thrust/iterator/constant_iterator.h>
...
// create iterators
thrust::constant_iterator<int> first(10);
thrust::constant_iterator<int> last = first + 3;

first[0]   // returns 10
first[1]   // returns 10
first[100] // returns 10

// sum of [first, last)
thrust::reduce(first, last);   // returns 30 (i.e. 3 * 10)

  每当需要恒定值的输入序列时，这constant_iterator是一种方便有效的解决方案。

transform_iterator

  在算法部分，我们讨论了内核融合，即将单独的算法组合在一起，transform并将其简化为单个transform_reduce操作。将transform_iterator允许我们采用相同的技术，即使我们没有特殊transform_xxx的算法的版本。这个例子展示了另一种融合转换与减少的方法，这次只用了简单的简化应用于a transform_iterator。

#include <thrust/iterator/transform_iterator.h>
// initialize vector
thrust::device_vector<int> vec(3);
vec[0] = 10; vec[1] = 20; vec[2] = 30;

// create iterator (type omitted)
... first = thrust::make_transform_iterator(vec.begin(), negate<int>());
... last  = thrust::make_transform_iterator(vec.end(),   negate<int>());

first[0]   // returns -10
first[1]   // returns -20
first[2]   // returns -30

// sum of [first, last)
thrust::reduce(first, last);   // returns -60 (i.e. -10 + -20 + -30)

  注意，为简单起见，我们首先省略了迭代器的类型。一个缺点transform_iterator是指定迭代器的完整类型可能很麻烦，这可能非常冗长。出于这个原因，通常的做法是简单地将调用放入make_transform_iterator被调用的算法的参数中。例如，

// sum of [first, last)
thrust::reduce(thrust::make_transform_iterator(vec.begin(), negate<int>()),
               thrust::make_transform_iterator(vec.end(),   negate<int>()));

  允许我们避免创建存储first和变量last。

permutation_iterator

  在上一节中，我们展示了如何transform_iterator使用另一种算法融合转换以避免不必要的内存操作。它permutation_iterator是类似的：它允许我们使用Thrust算法或甚至其他花哨的迭代器融合聚集和分散操作。以下示例显示如何将收集操作与减少融合。

#include <thrust/iterator/permutation_iterator.h>
...
// gather locations
thrust::device_vector<int> map(4);
map[0] = 3;
map[1] = 1;
map[2] = 0;
map[3] = 5;

// array to gather from
thrust::device_vector<int> source(6);
source[0] = 10;
source[1] = 20;
source[2] = 30;
source[3] = 40;
source[4] = 50;
source[5] = 60;

// fuse gather with reduction: 
//   sum = source[map[0]] + source[map[1]] + ...
int sum = thrust::reduce(thrust::make_permutation_iterator(source.begin(), map.begin()),
                         thrust::make_permutation_iterator(source.begin(), map.end()));

  这里我们使用该make_permutation_iterator函数来简化构造permutation_iterators。第一个参数make_permutation_iterator是收集操作的源数组，第二个参数是映射索引列表。请注意，我们source.begin()在两种情况下都传入第一个参数，但改变第二个参数以定义序列的开头和结尾。
  当a permutation_iterator用作函数的输出序列时，它等效于将散射操作融合到算法中。通常permutation_iterator允许您对序列中的一组特定值进行操作，而不是整个序列。

zip_iterator

  继续阅读，我们已经保存了最好的迭代器！这zip_iterator是一个非常有用的小工具：它需要多个输入序列并产生一系列元组。在这个例子中，我们将一系列序列int和一系列序列“压缩” char成一个序列，tuple<int,char>并tuple用最大值计算。

#include <thrust/iterator/zip_iterator.h>
...
// initialize vectors
thrust::device_vector<int>  A(3);
thrust::device_vector<char> B(3);
A[0] = 10;  A[1] = 20;  A[2] = 30;
B[0] = 'x'; B[1] = 'y'; B[2] = 'z';

// create iterator (type omitted)
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]   // returns tuple(10, 'x')
first[1]   // returns tuple(20, 'y')
first[2]   // returns tuple(30, 'z')

// maximum of [first, last)
thrust::maximum< tuple<int,char> > binary_op;
thrust::tuple<int,char> init = first[0];
thrust::reduce(first, last, init, binary_op); // returns tuple(30, 'z')

  使得zip_iterator如此有用的是大多数算法接受一个或偶尔两个输入序列。这zip_iterator允许我们将许多独立序列组合成单个元组序列，这可以通过一组广泛的算法进行处理。
  请参阅arbitrary_transformation示例，了解如何使用zip_iterator和实现三元变换for_each。此示例的简单扩展将允许您计算具有多个输出序列的转换。
  除了方便之外，还zip_iterator允许我们更有效地实施计划。例如，将3d点存储为float3CUDA中的数组通常是一个坏主意，因为数组访问未正确合并。随着zip_iterator我们可以在三个坐标存储在三个单独的数组，确实允许合并内存访问。在这种情况下，我们使用zip_iterator创建一个3d矢量的虚拟数组，我们可以将其输入Thrust算法。dot_products_with_zip有关其他详细信息，请参阅示例。

Additional Details

  本指南仅涉及Thrust可以做的事情。以下资源可以帮助您学习如何使用Thrust做更多事情或在出现问题时提供帮助。
  Thrust API的综合文档
  常见问题清单
  示例程序的集合
  我们强烈建议用户订阅推送用户邮件列表。邮件列表是寻求Thrust开发人员和其他Thrust用户帮助的好地方。

reference

https://github.com/thrust/thrust
https://github.com/thrust/thrust/wiki/Quick-Start-Guide
http://thrust.github.com