二战结束后,考虑到二战为人类带来的巨大灾难,爱因斯坦与特斯拉联手研发了一台时空穿梭机,并回到了1924年,除掉了由于啤酒馆政变入狱的希特勒,纳粹德国不复存在,但这却将欧洲拖入了新的血雨腥风,使得苏联统治了整个欧洲。不久,斯大林被Nod兄弟会派来的女刺客暗杀……
斯大林死后,Nod兄弟会潜伏在苏联的皮囊之下,直到1995年,一场神秘的陨石雨降临地球,带来了一种叫做“泰伯利亚”的危险有毒,能迅速增殖,且有很高的经济价值的外星矿物,它的蔓延将使地球陷入到核冬天。联合国主导的全球防卫组织努力消灭“泰伯利亚”,而NOD兄弟会的领袖Kane则试图利用“泰伯利亚”获取经济利益,并实现统治世界的野心。
最终NOD兄弟会和全球防卫组织谁胜利了呢?这取决于《命令与征服》和《红色警戒》的游戏玩家。
在游戏中,Kane的一句话深入人心:
He who commands the future conquers the past。He who conquers the past commands the future.
方老师作为《命令与征服》的玩家,把这句话抄在了几本书的扉页上,并且在《算法导论》的扉页上改成了:
He who divides the future conquers the past。He who conquers the past divides the future。
这是因为,divide and conquer (分治)算法是非常有价值的算法。
所谓的分治,就是将一个大而复杂的问题,拆分为小而容易解决的问题。在上一期中提到的矩阵乘法也可以使用分治算法:
如果每个Ann和Bnn均为一个4x4的矩阵,实际上就可以把8x8的矩阵乘法转换为做8次4x4的矩阵乘法,从而利用Volta的Tensor Core实现。
我们知道,需要做这些运算的数据都在CPU挂载的内存里面,那么,我们应当如何让GPU计算它们呢?
上图是一台典型的Intel x86 v7服务器的架构框图,GPU通过PCI-E总线与CPU相连,GPU也可以利用PCI-E的MSI中断和DMA机制从系统内存中读取数据。
GPU运算的流程如上图:
- CPU告诉GPU,把系统内存中的数据复制到GPU内存;
- CPU把GPU指令传给GPU;
- GPU的各个计算单元并行执行运算;
- GPU将计算结果复制回系统内存;
那么,工程师们在编写GPU运算程序的时候,是不是要处理DMA和中断呢?
答案当然是否定的。
2008年,Nvidia推出了CUDA(Compute Unified Device Architecture,统一计算架构),贯穿了编译器、运行时库、驱动和硬件层。开发者只要基于CUDA的规范编写程序就可以实现用C语言编写并行计算的算法,并让GPU执行该算法,而无需关心CPU和GPU交互的细节。
让我们举一个栗子:
这是一段执行向量加法的代码:
代码语言:javascript复制 // Vecadd definition
void VecAdd(float* A, float* B, float* C, unsigned int N)
{
for (int i = 0; i<N; i )
{
C[i] = A[i] B[i];
}
}
int main()
{
...
// invocate vector add function for CPU
VecAdd (A, B, C, N);
...
}
第2-8行定义了向量加法函数,将2个N维向量A和B相加得到C。有C语言基础的同学都可以看出,这是一个时间复杂度为O(N)的算法。
调用CUDA的算法代码如下:
代码语言:javascript复制//Kernel definition
__global__ void VecAdd(float* A, float* B, float* C)
{
int i = threadIdx.x;
C[i] = A[i] B[i];
}
int main()
{
...
// Kernel invocation with N threads
VecAdd<<<1, N>>>(A, B, C);
...
}
这段代码和以上代码最大的区别在于,向量加函数VecAdd调用了CUDA库,利用GPU并行计算向量运算,也就是所谓的“CUDA kernel”。经编译器进行编译连接后,生成的代码会调用CUDA库框架,实现这些功能:
- CPU告诉GPU,把系统内存中的数据复制到GPU内存;
- CPU把GPU指令传给GPU;
- GPU的各个计算单元并行执行运算;
- GPU将计算结果复制回系统内存;
类似地,我们还可以利用CUDA实现向量乘法、矩阵乘法等算法。
这样一来,工程师们就可以方便地利用GPU来进行并行运算了。