持续创作，加速成长！这是我参与「掘金日新计划 · 10 月更文挑战」的第20天，点击查看活动详情
上一篇文章是对核函数的开篇，介绍了如何编写简单的核函数，只不过我们使用了pycuda提供的函数模板，有时候我们可能为了简化操作，让我们的核函数能够适应很多任务，而不局限于某种任务，因此我们需要设计自己的核函数，让它方便完成我们特定的几种任务，而并非单一任务。所以SourceModule函数应运而生，它的出现就是为了能够让设计者可以方便的设计自己的核函数，而同时又能够得到CUDA的正确编译。
使用SourceModule编写的内核函数不必遵循pycuda提供的严格的设计模式，这样可以使得核函数的功能得到很大的拓展，但是相应的编程的复杂性也会大大的提高。
SourceModule函数会将原来的内联CUDA C代码进行编译，形成可以从python启动的内核函数，实际上，这里面的流程是：SourceModule会先将这些代码进行编译，编译成可以跟CUDA直接交互的模块，可以理解为DLL文件吧，只是举个例子方便理解，然后就可以通过pycuda模块对这些编译好的模块进行调用了。
下面通过一个简单的向量与标量的乘法来看一下使用SourceModule是如何编写内核函数的。
import pycuda.autoinit
import pycuda.driver as drv
import numpy as np
from pycuda import gpuarray
from pycuda.compiler import SourceModule
# 核函数
kernal_function = SourceModule(
"""
__global__ void vector_multiple(float *vecin, float *vecout, float delta)
{
int i = threadIdx.x;
vecout[i] = `delta * vecin[i]`;
}
"""
)
# 取出编译好的内核函数
kernal_function_togpu = kernal_function.get_function("vector_multiple")
vecin = np.random.rand(512).astype(np.float32)
vecin_gpu = gpuarray.to_gpu(vecin)
# 在GPU上新建一个与vecin_gpu同样大小的空向量
vecout = gpuarray.empty_like(vecin_gpu)
kernal_function_togpu(vecin_gpu, vecout, np.float32(3), block=(512, 1, 1), grid=(1, 1, 1))
print("source vector:\n", vecin_gpu.get())
print("result output:\n{}".format(vecout.get()))
先不看整体代码，就看SourceModule这一块，我把他们的作用逐行介绍一下
__global__：作为一组关键字参数，它的作用是向编译器进行声明，这是一个内核函数。
void vector_multiple(float *vecin, float *vecout, float delta)：这个含明显，就参照C语言来解释，前面的void就是函数vector_multiple的类型，表明此函数没有返回值，函数vector_multiple里面有三个参数，其中第一个是一个float类型的指针变量，在此处的含义是一个存放浮点类型输入数据的指针，第二个是存放浮点类型的输出数据的指针，第三个是一个标量，就单纯的一个float变量。
int i = threadIdx.x;：借鉴前面的理解，在ElementwiseKernal中函数是通过i值在GPU上实现自动并行的，而且这个值由pycuda控制，这个地方可以理解为，每个线程都有独立的标识，这个标识是由threadIdx设置的，所以对于每个线程的索引就是依靠i进行的。
vecout[i] = delta * vecin[i];这句，不能纯使用C语言来看，就结合python中使用numpy时向量的乘法运算来看，非常好理解，就是一个标量与一个向量的乘积。
kernal_function_togpu = kernal_function.get_function("vector_multiple")这一句，就是从SourceModule编译的CUDA模块中取出已经编译好的内核函数的引用。可能这句话有点绕，如果实在记不住，你就理解为实例化吧。
kernal_function_togpu(vecin_gpu, vecout, np.float32(3), block=(512, 1, 1), grid=(1, 1, 1))：参数block的格式是一个长度为3的元组，从几何的角度来说便是格点长方体各个边的格点数量，格点长方体中的每一个格点便是一个线程。grid与block是类似的，只不过每一个格点不是一个线程，而是一个完整的block！举个例子，有一所小学，每个班级都有3列4排桌椅共可容纳12个学生。我们有一个6层的教学楼（一层一个年级，1-6楼对应一至六年级），每层8间教室（1-8班），共计48间教室（一个班级占据一个完整的教室）。一个线程就是一个学生，一个block就是一间教室，一个grid就是一层教学楼，GPU就是这栋教学楼
具体的对应关系看下图。
最后就是，代码的运行截图
