从STL到THRUST，CUDA代码封装诀窍，让CUDA代码像C++一样优美

liuian 2024-12-23 12:11 66 浏览

cuda代码虽然可以使用计算设备达到超高并行化，超高吞吐的效果，但是写起来大部分面向过程，优美性和灵活性都有所欠缺，这就促使我们思考，能不能像C封装成C++一样，做一些简单封装，让二者的融合度更高，让CUDA的编程更友好？

thrust应运而生。它支持host和device两种vector，同时提供了类似STL的五大功能：

变换 (transformation)
规约 (reduction)
前缀和 (prefix sum)
排序 (sorting), 搜索(searching)
选择性复制、替换、移除、分区等重排(reordering) 操作

在这里，我们将分析几种关键功能实现和设计思路，方便理解和拓展thrust相关的开发。

使用中间层将C++类, CUDA代码分离

cuda中kernel的使用大致如下：

// 定义
__global__ void functionA_kernel_step(args) {}
__global__ void functionA_kernel_step2(args) {}
//调用:
functionA_kernel_step<<<1, 5>>>(args);
functionA_kernel_step2<<<1, 5>>>(args);

这样不仅繁琐，而且不易理解。C++中广泛采用的面向对象，一般会将函数描述成某个class的成员，类似：

class Master {
public:
Master();
~Master();
void CallMonster();
void DismissMonsters();
};
class Monster {
public:
Monster();
~Monster();
void Attack(Other something);
};

假设attack这里是核函数的形式调用，我们希望C++代码尽量整洁。

这点倒是易于攻破，只需要将kernel的调用封装到另一个文件里，C++中调用对应的功能即可。

// file: action.cu
__global__ void monster_attack() {
}
void CallAttack() {
kernel<<<3, 3>>>monster_attack();
}

// file: action.h
void CallAttack();

// file: monster.cpp
#include "action.h"
void Monster::Attack(Other something) {
CallAttack();
}

容器内存复制功能封装

这样组织文件可以让Cpp部分的代码独立一些，便于设计。kernel的实现并没有多少优化。我们还需要进一步实现一些必须的“轮子”。

C++中最广泛使用的数据结构是vector和map(unordered_map), map的设计过于复杂，这里不多加描述。如果是vector，排除realloc的部分，其实都是访问一块连续内存。

暂时不考虑iterator的封装的话，我们也可以设计一个简洁的，能直接从C++的vector中复制对应Vector的类。

class MasterIntVector {
public:
MasterIntVector(const std::vector<int> &array) {
n = array.size();
cudaMalloc((void**)(a), n * sizeof(int));
cudaMemcpy((void**)(a), array.data(), n * sizeof(int), cudaHostToDevice, 0);
}
MasterIntVector(int len) {
n = len;
cudaMalloc((void**)(a), n * sizeof(int));
cudaMemcpy((void**)(a), array.data(), n * sizeof(int), cudaHostToDevice, 0);
}
~MasterIntVector() {cudaFree(a);}
const int* begin() {
return a;
}
const int* end() {
return a + n;
}
private:
int *a;
}

自定义函数类型的封装和传递

有了这个简陋的封装类，我们就可以使用类似STL里的for_each功能，例如，对每个元素增加42。

// for_each.cu
__global__ void for_each_kernel(int *data, int n, func_type func) {
for (int i = 0; i < n; i++) func(data[i]);
}
void for_each(int *begin, int *end, func_type func) {
int n = end - begin;
for_each_kernel<<<1, end - begin>>>(begin, n, func);
}

// for_each.c++
typedef int(*func_type)(int);
__device__ __host__ int Add42(int x) {
return x+42;
}
for (int i=0; i<n; i++) {
cpu_monster.push_back(i); //创建一个顺序的vector
}
MasterIntVector a {cpu_monster};
for_each(a.begin(), a.end(), Add42); // WRONG!!! HOW????

这里就引入第三个需要处理的封装问题：如何在C++代码中向核函数传入对应的自定义函数；

一般在host端，能获得的函数地址都是host端的。直接将Add42这个地址传入，在device中并不能执行。__device__和__host__宏定义的意义在于，在kernel函数和host端，都能找到这个符号对应的地址，无法负责两个地址间的访问转换。

不过，有另一个使得函数内容能够拷贝到device的方式：使用struct定义functor。