CUDA从入门到放弃（十四）：CUDA Thrust库

码农世界 2024-06-19 后端 98 次浏览 0个评论

Thrust 是一个基于标准模板库（STL）的 C++ 模板库，专为 CUDA 设计，旨在简化高性能并行应用的开发。它提供了一系列数据并行原语，如扫描、排序和归约，可组合实现复杂算法。通过高级抽象描述计算，Thrust 能自动选择最优实现，适用于 CUDA 应用的快速原型设计和生产环境，提高程序员生产率和性能。

1 安装

安装 CUDA Toolkit 会将 Thrust 的头文件复制到您系统的标准 CUDA 包含目录中。由于 Thrust 是一个由头文件组成的模板库，因此无需进行其他安装步骤即可开始使用 Thrust。

2 Vectors

Thrust 提供了 host_vector 和 device_vector 两种向量容器，分别用于主机和 GPU 设备内存。它们类似于 C++ STL 中的 std::vector，是泛型容器，可动态调整大小。使用 = 运算符可以轻松地复制容器内容。device_vector 的元素可以通过标准括号访问，但应谨慎使用，因为每次访问都涉及内存复制。

#include 
#include 
#include 
int main(void)
{
    // H has storage for 4 integers
    thrust::host_vector 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 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 deleted when the function returns
    return 0;
}

Thrust 还提供了初始化向量元素和复制特定值集的方法。

#include 
#include 
#include 
#include 
#include 
#include 
int main(void)
{
    // initialize all ten integers of a device_vector to 1
    thrust::device_vector 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 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;
}

2-1 Thrust 命名空间

Thrust 命名空间允许我们调用特定的函数或类，如 thrust::host_vector 和 thrust::copy，从而避免与其他库中的函数或类名称冲突。

2-2 迭代器与静态分发 Iterators and Static Dispatching

在 Thrust 中，迭代器类似于指针，用于访问容器中的元素。Thrust 函数通过检查迭代器的类型，自动决定使用主机还是设备实现，这个过程称为静态分发，它在编译时确定，没有运行时开销。

原始指针和device_ptr 互相转换：

如果使用原始指针作为参数，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 dev_ptr(raw_ptr);
// use device_ptr in thrust algorithms
thrust::fill(dev_ptr, dev_ptr + N, (int) 0);

从 device_ptr 中提取原始指针：

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

STL 和互相转换：

#include 
#include 
#include 
#include 
int main(void)
{
    // create an STL list with 4 values
    std::list 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 D(stl_list.begin(), stl_list.end());
    // copy a device_vector into an STL vector
    std::vector stl_vector(D.size());
    thrust::copy(D.begin(), D.end(), stl_vector.begin());
    return 0;
}

3 Algorithms

Thrust 提供众多常见并行算法，其中许多与 STL 中的算法相对应，并采用相同名称。这些算法都有主机和设备的实现，根据使用的迭代器类型自动选择执行路径。除了 thrust::copy 外，Thrust 算法的迭代器参数应全部位于同一位置（主机或设备），否则编译器会报错。

3-1 Transformations

以下源代码演示了几种转换算法。

#include 
#include 
#include 
#include 
#include 
#include 
#include 
#include 
int main(void)
{
    // allocate three device_vectors with 10 elements
    thrust::device_vector X(10);
    thrust::device_vector Y(10);
    thrust::device_vector 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());
    // 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());
    // 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(std::cout, "\n"));
    return 0;
}

3-2 Reductions

归约算法使用二元操作将输入序列减少到单个值。

例如，通过使用加法操作对数字数组进行归约，可以得到数组的和。使用 thrust::reduce 实现数组的和如下：

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

通过使用接受两个输入并返回最大值的操作符进行归约，可以得到数组的最大值。

 int max_value = thrust::reduce(data.begin(), data.end(),  
                                    data[0], // 初始值  
                                    thrust::maximum()); // 自定义操作符

thrust::count 返回给定序列中特定值的实例数量。

int result = thrust::count(vec.begin(), vec.end(), 1);

3-3 Prefix-Sums

并行前缀和，或扫描操作，是许多并行算法（如流压缩和基数排序）中的重要组成部分。

以下代码演示了使用默认加法运算符的 inclusive_scan 操作：

#include 
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]。

exclusive scan 类似，但向右移动了一个位置：

#include 
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]。

Thrust 还提供了 transform_inclusive_scan 和 transform_exclusive_scan 函数，它们在执行扫描之前对输入序列应用一元函数。

3-4 Reordering

Thrust 通过以下算法提供了对分区和流压缩的支持：

copy_if：复制满足特定断言条件的元素到一个新的容器

#include   
#include   
#include  // for thrust::is_even  
#include   
  
int main(void)  
{  
    // 初始化设备向量  
    thrust::device_vector data = {1, 2, 3, 4, 5, 6, 7, 8, 9, 10};  
      
    // 创建一个空的设备向量用于存储结果  
    thrust::device_vector even_numbers(data.size());  
      
    // 使用 copy_if 复制偶数  
    // 注意：copy_if 的结果迭代器是输出范围的开始迭代器  
    thrust::copy_if(data.begin(), data.end(),  
                    even_numbers.begin(),  
                    thrust::is_even()); // 断言条件：检查是否为偶数  
      
    // 获取实际复制的元素数量  
    size_t even_count = thrust::distance(even_numbers.begin(),  
                                          thrust::remove_if(even_numbers.begin(), even_numbers.end(), thrust::not1(thrust::is_even())).end());  
      
    // 调整 even_numbers 的大小以匹配实际元素数量  
    even_numbers.resize(even_count);  
      
    // 输出结果  
    std::cout << "Even numbers: ";  
    thrust::copy(even_numbers.begin(), even_numbers.end(), std::ostream_iterator(std::cout, " "));  
    std::cout << std::endl;  
      
    return 0;  
}

partition：partition 函数会根据提供的断言条件重新排列输入序列中的元素，使得满足条件的元素出现在不满足条件的元素之前。

#include   
#include   
#include  // for thrust::less  
#include   
  
int main(void)  
{  
    // 初始化设备向量  
    thrust::device_vector data = {1, 3, 2, 5, 4, 6, 7, 8, 9, 10};  
      
    // 使用 partition 重新排列数据，使得小于 5 的数在前  
    thrust::partition(data.begin(), data.end(), thrust::less(5));  
      
    // 输出结果  
    std::cout << "Partitioned data: ";  
    thrust::copy(data.begin(), data.end(), std::ostream_iterator(std::cout, " "));  
    std::cout << std::endl;  
      
    return 0;  
}

remove 和 remove_if：remove函数会移除给定序列中等于某个特定值的所有元素，而remove_if函数则会根据提供的断言条件移除元素。

#include   
#include   
#include   
#include   
  
int main(void)  
{  
    // 初始化设备向量  
    thrust::device_vector data = {1, 2, 3, 2, 4, 2, 5, 6};  
      
    // 使用 remove 移除所有值为 2 的元素  
    thrust::device_vector::iterator new_end = thrust::remove(data.begin(), data.end(), 2);  
	// 使用 remove_if 移除所有大于 5 的元素  
    // thrust::device_vector::iterator new_end = thrust::remove_if(data.begin(), data.end(), thrust::greater(5));
      
    // 调整数据向量的大小以匹配新的末尾迭代器  
    data.resize(thrust::distance(data.begin(), new_end));  
      
    // 输出结果  
    std::cout << "Data after removing 2s: ";  
    thrust::copy(data.begin(), data.end(), std::ostream_iterator(std::cout, " "));  
    std::cout << std::endl;  
      
    return 0;  
}

unique：unique函数会移除序列中连续重复的元素，使得每个元素只出现一次。

#include   
#include   
#include   
#include   
  
int main(void)  
{  
    // 初始化设备向量  
    thrust::device_vector data = {1, 2, 2, 3, 4, 4, 4, 5, 6, 6};  
      
    // 使用 unique 移除连续重复的元素  
    thrust::device_vector::iterator new_end = thrust::unique(data.begin(), data.end());  
      
    // 调整数据向量的大小以匹配新的末尾迭代器  
    data.resize(thrust::distance(data.begin(), new_end));  
      
    // 输出结果  
    std::cout << "Data after removing consecutive duplicates: ";  
    thrust::copy(data.begin(), data.end(), std::ostream_iterator(std::cout, " "));  
    std::cout << std::endl;  
      
    return 0;  
}

3-5 Sorting

Thrust 提供多种函数，可按给定条件排序或重新排列数据。其中，thrust::sort 和 thrust::stable_sort 与 STL 中的对应函数相似。

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

4 Fancy Iterators

4-1 constant_iterator

constant_iterator 是一种简单的迭代器，它在每次访问时返回相同的值。

#include 
...
// create iterators
thrust::constant_iterator first(10);
thrust::constant_iterator 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)

4-2 counting_iterator

counting_iterator 生成递增的序列。

#include 
...
// create iterators
thrust::counting_iterator first(10);
thrust::counting_iterator last = first + 3;
first[0]   // returns 10
first[1]   // returns 11
first[100] // returns 110
// sum of [first, last)
thrust::reduce(first, last);   // returns 33 (i.e. 10 + 11 + 12)

4-3 transform_iterator

transform_iterator 允许我们对序列中的每个元素应用转换。

#include 
// initialize vector
thrust::device_vector vec(3);
vec[0] = 10; vec[1] = 20; vec[2] = 30;
// create iterator (type omitted)
...
first = thrust::make_transform_iterator(vec.begin(), negate());
...
last  = thrust::make_transform_iterator(vec.end(),   negate());
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)
// 或者
// sum of [first, last)
thrust::reduce(thrust::make_transform_iterator(vec.begin(), negate()),
               thrust::make_transform_iterator(vec.end(),   negate()));

4-4 permutation_iterator

permutation_iterator 允许我们重新排列序列中元素的顺序。通过指定一个映射，我们可以按特定顺序访问元素：

#include 
...
// gather locations
thrust::device_vector map(4);
map[0] = 3;
map[1] = 1;
map[2] = 0;
map[3] = 5;
// array to gather from
thrust::device_vector 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()));

4-5 zip_iterator

zip_iterator 可以将多个序列合并为一个元组序列。这使得我们可以同时处理多个序列：

#include 
...
// initialize vectors
thrust::device_vector  A(3);
thrust::device_vector 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 > binary_op;
thrust::tuple init = first[0];
thrust::reduce(first, last, init, binary_op); // returns tuple(30, 'z')

参考资料

1 Thrust docs

2 Thrust: The C++ Parallel Algorithms Library

转载请注明来自码农世界，本文标题：《CUDA从入门到放弃（十四）：CUDA Thrust库》

码农世界 37815篇文章站点微博

每一天，每一秒，你所做的决定都会改变你的人生！

发表评论取消回复

评论列表（暂无评论，98人围观）参与讨论

码农世界管理员

CUDA从入门到放弃（十四）：CUDA Thrust库

1 安装

2 Vectors

2-1 Thrust 命名空间

2-2 迭代器与静态分发 Iterators and Static Dispatching

3 Algorithms

3-1 Transformations

3-2 Reductions

3-3 Prefix-Sums

3-4 Reordering

3-5 Sorting

4 Fancy Iterators

4-1 constant_iterator

4-2 counting_iterator

4-3 transform_iterator

4-4 permutation_iterator

4-5 zip_iterator

发表评论取消回复

还没有评论，来说两句吧...

文章目录

码农世界管理员

CUDA从入门到放弃（十四）：CUDA Thrust库

1 安装

2 Vectors

2-1 Thrust 命名空间

2-2 迭代器与静态分发 Iterators and Static Dispatching

3 Algorithms

3-1 Transformations

3-2 Reductions

3-3 Prefix-Sums

3-4 Reordering

3-5 Sorting

4 Fancy Iterators

4-1 constant_iterator

4-2 counting_iterator

4-3 transform_iterator

4-4 permutation_iterator

4-5 zip_iterator

Visual Studio 连接 MySQL 数据库 实现数据库的读写（C++）

[工业自动化-1]：PLC架构与工作原理

网络编程（六）TCP并发服务器

Vue3、Element Plus使用v-for循环el-form表单进行校验

Ceph: vdbench 测试ceph存储rbd块设备

linux如何部署前端项目和安装nginx

Windows OpenVPN的安装之服务器自动启动连接

服务器数据恢复—KVM虚拟机被误删除如何恢复虚拟磁盘文件？

发表评论取消回复

还没有评论，来说两句吧...

文章目录

Visual Studio 连接 MySQL 数据库实现数据库的读写（C++）