MultithreadCorner/thrust-multi-permutation-iterator

This repository adds functionality missing from the official version of Thrust. Multi_permutation_iterator is a general-purpose iterator, analogous to permutation_iterator that allows for tuple-based indexing into data -- think coalesced multi-dimensional arrays or stencil-based grid traversal. Without multi_permutation_iterator one must resort to redundantly applying zip_iterator to the same iterator perhaps with different offsets, resulting in excessive kernel argument size and overhead, and most importantly consumption of performance-critical GPU registers. Multi_permutation_iterator is completely general-purpose, providing functionality that zip_iterator+permutation_iterator only provides a crude workaround for. Another key missing piece of functionality is support for streaming tuple values via C++ streams.

Thrust: Code at the speed of light

Thrust is a parallel algorithms library which resembles the C++ Standard Template Library (STL). Thrust's high-level interface greatly enhances programmer productivity while enabling performance portability between GPUs and multicore CPUs. Interoperability with established technologies (such as CUDA, TBB, and OpenMP) facilitates integration with existing software. Develop high-performance applications rapidly with Thrust!

Examples

Thrust is best explained through examples. The following source code generates random numbers serially and then transfers them to a parallel device where they are sorted.

#include <thrust/host_vector.h>
#include <thrust/device_vector.h>
#include <thrust/generate.h>
#include <thrust/sort.h>
#include <thrust/copy.h>
#include <algorithm>
#include <cstdlib>

int main(void)
{
  // generate 32M random numbers serially
  thrust::host_vector<int> h_vec(32 << 20);
  std::generate(h_vec.begin(), h_vec.end(), rand);

  // transfer data to the device
  thrust::device_vector<int> d_vec = h_vec;

  // sort data on the device (846M keys per second on GeForce GTX 480)
  thrust::sort(d_vec.begin(), d_vec.end());

  // transfer data back to host
  thrust::copy(d_vec.begin(), d_vec.end(), h_vec.begin());

  return 0;
}

This code sample computes the sum of 100 random numbers in parallel:

#include <thrust/host_vector.h>
#include <thrust/device_vector.h>
#include <thrust/generate.h>
#include <thrust/reduce.h>
#include <thrust/functional.h>
#include <algorithm>
#include <cstdlib>

int main(void)
{
  // generate random data serially
  thrust::host_vector<int> h_vec(100);
  std:generate(h_vec.begin(), h_vec.end(), rand);

  // transfer to device and compute sum
  thrust::device_vector<int> d_vec = h_vec;
  int x = thrust::reduce(d_vec.begin(), d_vec.end(), 0, thrust::plus<int>());
  return 0;
}

Refer to the Quick Start Guide page for further information and examples.

Contributors

The primary developers of Thrust are Jared Hoberock and Nathan Bell.