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。

struct Add42 {

__device__ __host__ int operator()(int x) {

return x + 42;

}

}

调用的时候，需要创建一个Add42类型的实例，同时把func_type的定义修改下；

using func_type = Add42;

for_each(a.begin(), a.end(), Add42{});

这样，我们就实现了可拓展的for_each操作。