CUDA初学者-Thrust - Algorithms - Prefix Sums（3）

MQTXWD

已于 2022-05-16 08:55:19 修改

阅读量611

点赞数

分类专栏： cuda小白文章标签： c++

于 2022-05-13 09:26:31 首次发布

本文链接：https://blog.csdn.net/u011732139/article/details/124717147

版权

cuda小白专栏收录该内容

14 篇文章 1 订阅

订阅专栏

本人 CUDA小白一枚，要是有什么不对，还望各位大佬指点。
本文及后面的几篇将分别从几个方面来大概阐述一下Thrust的一些接口。原来的网址在这里。

1.Algorithms

1.3 Prefix Sums
1.3.1 inclusive_scan

template <typename DerivedPolicy, typename InputIterator, typename OutputIterator>
__host__ __device__ OutputIterator thrust::inclusive_scan(
	const thrust::detail::execution_policy_base< DerivedPolicy > & exec,
  	InputIterator first,
  	InputIterator last,
  	OutputIterator result
);

template <typename InputIterator, typename OutputIterator>
OutputIterator thrust::inclusive_scan(
  	InputIterator first,
  	InputIterator last,
  	OutputIterator result
);

template <typename DerivedPolicy, typename InputIterator, typename OutputIterator, typename AssociativeOperator>
__host__ __device__ OutputIterator thrust::inclusive_scan(
	const thrust::detail::execution_policy_base< DerivedPolicy > & exec,
  	InputIterator first,
  	InputIterator last,
  	OutputIterator result,
  	AssociativeOperator binary_op
);

template <typename InputIterator, typename OutputIterator, typename AssociativeOperator>
OutputIterator thrust::inclusive_scan(
  	InputIterator first,
  	InputIterator last,
  	OutputIterator result,
  	AssociativeOperator binary_op
);

将[first, last)中每个index对应的result的结果为[first, index]之间的数值进行操作，默认为求和，可以指定特殊的操作。
例子：

int data[6] = {1, 0, 2, 2, 1, 3};
thrust::inclusive_scan(thrust::host, data, data + 6, data);
// thrust::inclusive_scan(data, data + 6, data);
// data中的结果为{1, 1, 3, 5, 6, 9}

int data2[10] = {-5, 0, 2, -3, 2, 4, 0, -1, 2, 8};
thrust::maximum<int> binary_op;   // 取最大值
thrust::inclusive_scan(thrust::host, data2, data2+ 6, data2);
// thrust::inclusive_scan(thrust::host, data2, data2+ 6, data2);
// data2中的结果为{-5, 0, 2, 2, 2,4, 4, 4, 4, 8}

1.3.2 exclusive_scan

template <typename DerivedPolicy, typename InputIterator, typename OutputIterator>
__host__ __device__ OutputIterator thrust::exclusive_scan(
	const thrust::detail::execution_policy_base< DerivedPolicy > & exec,
  	InputIterator first,
  	InputIterator last,
  	OutputIterator result
);

template <typename InputIterator, typename OutputIterator>
OutputIterator thrust::exclusive_scan(
  	InputIterator first,
  	InputIterator last,
  	OutputIterator result
);

template <typename DerivedPolicy, typename InputIterator, typename OutputIterator, typename T>
__host__ __device__ OutputIterator thrust::exclusive_scan(
	const thrust::detail::execution_policy_base< DerivedPolicy > & exec,
  	InputIterator first,
  	InputIterator last,
  	OutputIterator result,
  	T init
);

template <typename InputIterator, typename OutputIterator, typename T>
OutputIterator thrust::exclusive_scan(
  	InputIterator first,
  	InputIterator last,
  	OutputIterator result,
  	T init
);

template <typename DerivedPolicy, typename InputIterator, typename OutputIterator, typename T, typename AssociativeOperator>
__host__ __device__ OutputIterator thrust::exclusive_scan(
	const thrust::detail::execution_policy_base< DerivedPolicy > & exec,
  	InputIterator first,
  	InputIterator last,
  	OutputIterator result,
  	T init,
  	AssociativeOperator binary_op
);

template <typename InputIterator, typename OutputIterator, typename T, typename AssociativeOperator>
OutputIterator thrust::exclusive_scan(
  	InputIterator first,
  	InputIterator last,
  	OutputIterator result,
  	T init,
  	AssociativeOperator binary_op
);

将[first, last)中每个index对应的result的结果为[first, index)之间的数值进行操作，默认为求和，可以指定特殊的操作，也可以指定初始化的值。与exclusive唯一的区别在于当前index是否包含在内。指定的初始化值主要是针对第一个index的时候，相当于在原来的数组前增加一个值。

int data[6] = {1, 0, 2, 2, 1, 3};
thrust::exclusive_scan(thrust::host, data, data + 6, data);
// thrust::exclusive_scan(data, data + 6, data);
// data输出 {0, 1, 1, 3, 5, 6}

int data2[6] = {1, 0, 2, 2, 1, 3};
thrust::exclusive_scan(thrust::host, data, data + 6, data, 4);
// thrust::exclusive_scan(data, data + 6, data, 4);
// data2 {4, 5, 5, 7, 9, 10}

int data3[10] = {-5, 0, 2, -3, 2, 4, 0, -1, 2, 8};
thrust::maximum<int> binary_op;
thrust::exclusive_scan(thrust::host, data, data + 10, data, 1, binary_op);
// thrust::exclusive_scan(data, data + 10, data, 1, binary_op);
// data3 {1, 1, 1, 2, 2, 2, 4, 4, 4, 4}

1.3.3 inclusive_scan_by_key

template <typename DerivedPolicy, typename InputIterator1, typename InputIterator2, typename OutputIterator>
__host__ __device__ OutputIterator thrust::inclusive_scan_by_key(
	const thrust::detail::execution_policy_base< DerivedPolicy > & exec,
 	InputIterator1 first1,
  	InputIterator1 last1,
  	InputIterator2 first2,
  	OutputIterator result
);

template <typename InputIterator1, typename InputIterator2, typename OutputIterator>
OutputIterator thrust::inclusive_scan_by_key(
 	InputIterator1 first1,
  	InputIterator1 last1,
  	InputIterator2 first2,
  	OutputIterator result
);

template <typename DerivedPolicy, typename InputIterator1, typename InputIterator2, typename OutputIterator, typename BinaryPredicate>
__host__ __device__ OutputIterator thrust::inclusive_scan_by_key(
	const thrust::detail::execution_policy_base< DerivedPolicy > & exec,
 	InputIterator1 first1,
  	InputIterator1 last1,
  	InputIterator2 first2,
  	OutputIterator result,
  	BinaryPredicate binary_pred
);

template <typename InputIterator1, typename InputIterator2, typename OutputIterator,  typename BinaryPredicate>
OutputIterator thrust::inclusive_scan_by_key(
 	InputIterator1 first1,
  	InputIterator1 last1,
  	InputIterator2 first2,
  	OutputIterator result,
  	BinaryPredicate binary_pred
);

template <typename DerivedPolicy, typename InputIterator1, typename InputIterator2, typename OutputIterator, typename BinaryPredicate, typename AssociativeOperator>
__host__ __device__ OutputIterator thrust::inclusive_scan_by_key(
	const thrust::detail::execution_policy_base< DerivedPolicy > & exec,
 	InputIterator1 first1,
  	InputIterator1 last1,
  	InputIterator2 first2,
  	OutputIterator result,
  	BinaryPredicate binary_pred,
  	AssociativeOperator binary_op
);

template <typename InputIterator1, typename InputIterator2, typename OutputIterator,  typename BinaryPredicate, typename AssociativeOperator>
OutputIterator thrust::inclusive_scan_by_key(
 	InputIterator1 first1,
  	InputIterator1 last1,
  	InputIterator2 first2,
  	OutputIterator result,
  	BinaryPredicate binary_pred,
  	AssociativeOperator binary_op
);

针对[first1, last1)中的数值如果出现连续相同的数值，称为一个segment。每一个segment中执行inclusive的操作。可以指定segment定义的方法，也可以指定同一个segment中需要完成的操作。
例子：

int A[10] = {1, 1, 1, 1, 1, 1, 1, 1, 1, 1};
int B[10] = {2, 1, 2, 2, 3, 4, 5, 5, 5, 5};
thrust::inclusive_scan_by_key(thrust::host, B, B + 10, A, A);
// thrust::inclusive_scan_by_key(B, B + 10, A, A);
// A的结果为 {1, 1, 1, 2, 1, 1, 1, 2, 3, 4}
// 如果 A[10] = {1, 0, 1, 0, 1, 0, 1, 0, 1, 0};
// 则对应的输出结果为 {1, 0, 1, 1, 1, 0, 1, 1, 2, 2}

int C[10] = {1, 0, 1, 0, 1, 0, 1, 0, 1, 0};
thrust::equal_to<int> binary_pred;
thrust::inclusive_scan_by_key(thrust::host, B, B + 10, C, C, binary_pred);
// thrust::inclusive_scan_by_key(B, B + 10, C, C, binary_pred);
// C的结果为{1, 0, 1, 1, 1, 0, 1, 1, 2, 2}

int D[10] = {1, 0, 1, 0, 1, 0, 1, 0, 1, 0};
thrust::plus<int> binary_op;
thrust::inclusive_scan_by_key(thrust::host, B, B + 10, C, C, binary_pred, binary_op);
// thrust::inclusive_scan_by_key(B, B + 10, C, C, binary_pred, binary_op);
// C的结果为{1, 0, 1, 1, 1, 0, 1, 1, 2, 2}

1.3.4 exclusive_scan_by_key

template <typename DerivedPolicy, typename InputIterator1, typename InputIterator2, typename OutputIterator>
__host__ __device__ OutputIterator thrust::exclusive_scan_by_key(
	const thrust::detail::execution_policy_base< DerivedPolicy > & exec,
 	InputIterator1 first1,
  	InputIterator1 last1,
  	InputIterator2 first2,
  	OutputIterator result
);

template <typename InputIterator1, typename InputIterator2, typename OutputIterator>
OutputIterator thrust::exclusive_scan_by_key(
 	InputIterator1 first1,
  	InputIterator1 last1,
  	InputIterator2 first2,
  	OutputIterator result
);

template <typename DerivedPolicy, typename InputIterator1, typename InputIterator2, typename OutputIterator, typename T>
__host__ __device__ OutputIterator thrust::exclusive_scan_by_key(
	const thrust::detail::execution_policy_base< DerivedPolicy > & exec,
 	InputIterator1 first1,
  	InputIterator1 last1,
  	InputIterator2 first2,
  	OutputIterator result,
  	T init
);

template <typename InputIterator1, typename InputIterator2, typename OutputIterator, typename T>
OutputIterator thrust::exclusive_scan_by_key(
 	InputIterator1 first1,
  	InputIterator1 last1,
  	InputIterator2 first2,
  	OutputIterator result,
  	T init
);

template <typename DerivedPolicy, typename InputIterator1, typename InputIterator2, typename OutputIterator, typename T, typename BinaryPredicate>
__host__ __device__ OutputIterator thrust::exclusive_scan_by_key(
	const thrust::detail::execution_policy_base< DerivedPolicy > & exec,
 	InputIterator1 first1,
  	InputIterator1 last1,
  	InputIterator2 first2,
  	OutputIterator result,
  	T init,
  	BinaryPredicate binary_pred
);

template <typename InputIterator1, typename InputIterator2, typename OutputIterator, typename T, typename BinaryPredicate>
OutputIterator thrust::exclusive_scan_by_key(
 	InputIterator1 first1,
  	InputIterator1 last1,
  	InputIterator2 first2,
  	OutputIterator result,
  	T init,
  	BinaryPredicate binary_pred
);

template <typename DerivedPolicy, typename InputIterator1, typename InputIterator2, typename OutputIterator, typename T, typename BinaryPredicate, typename AssociativeOperator>
__host__ __device__ OutputIterator thrust::exclusive_scan_by_key(
	const thrust::detail::execution_policy_base< DerivedPolicy > & exec,
 	InputIterator1 first1,
  	InputIterator1 last1,
  	InputIterator2 first2,
  	OutputIterator result,
  	T init,
  	BinaryPredicate binary_pred,
  	AssociativeOperator binary_op
);

template <typename InputIterator1, typename InputIterator2, typename OutputIterator, typename T, typename BinaryPredicate, typename AssociativeOperator>
OutputIterator thrust::exclusive_scan_by_key(
 	InputIterator1 first1,
  	InputIterator1 last1,
  	InputIterator2 first2,
  	OutputIterator result,
  	T init,
  	BinaryPredicate binary_pred,
  	AssociativeOperator binary_op
);

原理同inclusive_scan_by_key，同样可以指定segment的划分方法，同一个segment中需要执行的操作，以及同一个segment中初始的值。
例子：

int keys[10] = {0, 0, 0, 1, 1, 2, 3, 3, 3, 3};
int vals[10] = {1, 1, 1, 1, 1, 1, 1, 1, 1, 1};
thrust::exclusive_scan_by_key(thrust::host, key, key + 10, vals, vals);
// thrust::exclusive_scan_by_key(key, key + 10, vals, vals);
// vals的值为 {0, 1, 2, 0, 1, 0, 0, 1, 2, 3}

int vals2[10] = {1, 1, 1, 1, 1, 1, 1, 1, 1, 1};
int init = 5;
thrust::exclusive_scan_by_key(thrust::host, key, key + 10, vals2, vals2, init);
// thrust::exclusive_scan_by_key(key, key + 10, vals, vals, init);
// vals的值为 {5, 6, 7, 5, 6, 5, 5, 6, 7, 8}

int vals3[10] = {1, 1, 1, 1, 1, 1, 1, 1, 1, 1};
thrust::equal_to<int> binary_op;
thrust::exclusive_scan_by_key(thrust::host, key, key + 10, vals3, vals3, init, binary_op);
// thrust::exclusive_scan_by_key(key, key + 10, vals, vals, init, binary_op);
// vals的值为 {5, 6, 7, 5, 6, 5, 5, 6, 7, 8}

1.3.5 transform_inclusive_scan

template <typename DerivedPolicy, typename InputIterator, typename OutputIterator, typename UnaryFunction, typename AssociativeOperator>
__host__ __device__ OutputIterator thrust::transform_inclusive_scan(
	const thrust::detail::execution_policy_base< DerivedPolicy > & exec,
  	InputIterator first,
  	InputIterator last,
  	OutputIterator result,
  	UnaryFunction unary_op,
  	AssociativeOperator binary_op
);

template <typename InputIterator, typename OutputIterator, typename UnaryFunction, typename AssociativeOperator>
OutputIterator thrust::transform_inclusive_scan(
  	InputIterator first,
  	InputIterator last,
  	OutputIterator result,
  	UnaryFunction unary_op,
  	AssociativeOperator binary_op
);

在原先inclusive_scan的基础上增加了transform的操作。
例子：

int data[8] = {1, 0, 2, 2, 1, 3};
thrust::negate<int> unary_op;
thrust::plus<int> binary_op;
thrust::transform_inclusive_scan(thrust::host, data, data + 6, data, unary_op, binary_op);
// thrust::transform_inclusive_scan(data, data + 6, data, unary_op, binary_op);
// data  {-1, -1, -3, -5, -6, -9}

1.3.6 transform_exclusive_scan

template <typename DerivedPolicy, typename InputIterator, typename OutputIterator, typename UnaryFunction, typename T, typename AssociativeOperator>
__host__ __device__ OutputIterator thrust::transform_exclusive_scan(
	const thrust::detail::execution_policy_base< DerivedPolicy > & exec,
  	InputIterator first,
  	InputIterator last,
  	OutputIterator result,
  	UnaryFunction unary_op,
  	T init,
  	AssociativeOperator binary_op
);

template <typename InputIterator, typename OutputIterator, typename UnaryFunction, typename T, typename AssociativeOperator>
OutputIterator thrust::transform_exclusive_scan(
  	InputIterator first,
  	InputIterator last,
  	OutputIterator result,
  	UnaryFunction unary_op,
  	T init,
  	AssociativeOperator binary_op
);

在exclusive_scan的基础上增加了transform的操作。
例子：

int data[6] = {1, 0, 2, 2, 1, 3};
thrust::negate<int> unary_op;
thrust::plus<int> binary_op;

thrust::transform_exclusive_scan(thrust:host, data, data + 6， data, unary_op, 4, binary_op);
// thrust::transform_exclusive_scan(data, data + 6， data, unary_op, 4, binary_op);
// data的值为 {4, 3, 3, 1, -1, -2}