一起养成写作习惯！这是我参与「掘金日新计划 · 4 月更文挑战」的第9天，点击查看活动详情。
在开始之前我们先来介绍两个概念，理解好这两个概念有助于理解 cuda 并行编程模型。
kernel
kernel 这是一个 GPU 编程出现的概念，这里多说一句我们做将 SVM 从仅解决线性问题推广到解决非线性问题，提到核函数，核函数会作用到每个样本来增加其特性。
那么在 GPU 上，核函数也是对于数据可以进行并行处理的函数。每一个核函数可以都运行在一个线程上对数据进行进行操作。
grid
grid 是一个比较抽象的概念，是为了方便程序员程序设计，一组线程组织形式，也是为了方便访问，如下图
在图中，我们以网格(grid)形式来组织线程，每一个线程上运行一个应用程序(函数)就是一个 kernel 而且这些 kernel 是相同，且相互并不依赖可并行的程序。
接下来我们通过一个具体问题，来介绍如何改变我们思路，用 cuda 并行计算来解决这个问题，问题很简单，就是我们有下面这样一个 4x3 的二维矩阵(数组)我们要对数组每一个元素进行增加 1 的操作
arr = np.array([
[1,0,2,3],
[0,1,2,2],
[0,2,3,2]
])
当然我们可以例如 numpy 的 arr + 1 操作来实现，这里我们为了解释 cpu 是如何运算将问题做的更普遍一些，就是我们通过对矩阵行列遍历来访问每一个元素，然后对每一个元素进行操作来实现。
def inc_one(A):
for row in range(A.shape[0]):
for col in range(A.shape[1]):
A[row,col] = A[row,col] + 1
return A
而在我们需要考虑是如何将数组每一个元素放置到 grid 上每一个 kernel 上，也就是我们要做的是定义 kernel 对每一个元素操作函数，然后如何将元素对一个 grid 一个线程上
@cuda.jit
def inc_one_kernel(A):
row,col = cuda.grid(2)
if row < A.shape[0] and col < A.shape[1]:
A[row,col] += 1
grid 的形状并不一定与要操作矩阵具有相同形状，例如图中，是 4x4 grid 来操作 4x3。
这里简单解释一下上面代码，我们用注解 cuda.jit 方式来修饰一个函数，这个函数中我们可以调用一些 cuda API 来利用 GPU 资源。搜索我们需要得到 cuda 每一个线程，我们可以调用 cuda.grid 传入 2 得到一个网格的线程，然后返回 row，col 我们可以通过 row col 来定位到一个线程。有两点需要注意，第一点就是 cuda 函数通常不会返回值，也就是意味着我们对矩阵每一个元素的操作是 in plae 修改原有数据，随意如果不想修改原有的数据，我们在传入函数前就需要对数据进行 copy 一份。第二点我们需要注意我们访问 thread 的索引在矩阵 A 是否存在。
实际上在这里还存在一层结构，就是 block 结构，从上图不难看出我们 grid 划分为 block 然后 block 又被划分为一个一个 thread ，在 thread 上运行 kernel 程序。
