去年AIGC大火,程序员都把注意力放在了最上层,而忽略了提供算力的最底层:GPU。
不过这也正常,就像很少人直接针对CPU编程一样,直接针对GPU编程的人也不多。
但是了解一下GPU编程,绝对大有好处。
今天先聊聊GPU编程,然后再聊聊一个CUDA这个新的生态系统,对编程细节不感兴趣的可以直接拉到最后。
对了,文末还有免费送书的福利。
CPU vs GPU
图片
CPU的设计目标是“尽可能地降低延时”
(1) 强大的ALU(算术逻辑单元),可以在很少的时钟周期内完成算术运算。
(2) 巨大的Cache:加快指令和数据的存取速度
(3) 复杂的逻辑控制:当程序员有多个分支,它可以通过分支预测来降低延时。
GPU的目标是:“尽可能地实现大吞吐量”
(1) ALU 简单,但是超级多
(2) Cache很小
(3) 逻辑控制简单。
如果把GPU的单个核心比作小学生,那一个CPU的核心就是老教授。
如果要做微积分,几千个小学生也比如上老教授。
但是,如果只是100以内的加减法,几千个小学生同时做(并行计算),那效率肯定要比老教授高。
老教授处理复杂任务的能力是碾压小学生的,但是对于没有那么复杂的任务,还是顶不住人多。
把串行改成并行
我们用一个例子来展示一下:
int a[] = {1,2,3,4,5,6,8,9,10};
int b[] = {11,12,13,14,15,16,17,18,19,20};
int c[10];
int main() {
int N = 10; // Number of elements
for (int i = 0; i < N; i++) {
c[i] = a[i] + b[i];
}
return 0;
}
这段简单的代码大家都能看懂,CPU在执行时会做一个循环,然后把两个数组对应的元素进行相加,结果存到数组c中。
由于是顺序处理的,如果数组非常大,就会比较耗时。
如何把它改成并行计算呢?
数组中有10个元素,我们可以创建10个线程,把每个线程扔到一个GPU核心中去运行。
图片
程序员该怎么写代码,来表达这个想法呢?
CUDA
英伟达的CUDA是一个并行计算平台,可以让程序员可以通过C、C++等语言在GPU上并行执行代码。
图片
在CUDA中,把CPU所在的部分叫做Host,GPU称为Device,它们之间通过总线相连。
图片
对于之前的例子,CUDA代码是这样的:
__global__ void vectorAdd(int* a, int* b, int* c){
int i = threadIdx.x;
c[i] = a[i] + b[i];
return;
}
估计大部分小伙伴都能猜出来这段代码的含义。
a,b分别是两个要想加的数组,c用来保存结果。
__global__应该是个指示符,表示这段代码是个“内核函数”,要被放到GPU上来执行。
threadIdx是个什么东西?
似乎是个线程的索引,找到这个线程的index以后,取出a,b中index对应的值,加起来放到c中。例如index是0,那就取出a[0],b[0]加起来,放到c[0]中,这就实现了我们之前的想法。
值得注意的是,这里的a,b,c不是Host的内存,而是Device(GPU)的内存,所以我们得把原始的数据复制到GPU中。
1. 先在GPU中分配内存
int* cudaA = 0;
int* cudaB = 0;
int* cudaC = 0;
// 使用cudaMalloc在GPU中分配内存
cudaMalloc(&cudaA,sizeof(a));
cudaMalloc(&cudaB,sizeof(b));
cudaMalloc(&cudaC,sizeof(c));
2.然后把原始数据从Host复制到Device(即GPU)中
//注意第4个参数,是从Host 到 Device
cudaMemcpy(cudaA, a, sizeof(a), cudaMemcpyHostToDevice);
cudaMemcpy(cudaB, b, sizeof(b), cudaMemcpyHostToDevice);
3. 调用内核函数
vectorAdd <<<1, sizeof(a) / sizeof(a[0])>>> (cudaA, cudaB, cudaC);
调用vectorAdd的时候,被<<< >>>包围起来的部分是配置参数,这里指定了一组10个线程(数组长度为10)。
这10个线程会被放到10个GPU核心中去执行,他们的索引是从0到9。
所以在vectorAdd函数中可以通过threadIdx.x引用到当前线程的索引,例如9 , 那就知道当前线程要做的事情:把a[9]和b[9]加起来,放到c[9]中。
这样10个GPU核心就是同时执行10次加法,速度飞快。
4. 把结果复制回Host
// 注意第4个参数,是从Device 到 Host
cudaMemcpy(c, cudaC, sizeof(c), cudaMemcpyDeviceToHost);
小伙伴们肯定已经意识到了,这里边有个核心的概念:Thread(线程),每个线程都会被映射到一个GPU核心去执行。
图片
多个Thread可以组成一个块(Block),被映射到多个核心
图片
多个Block又形成一个Grid,被映射到整个CPU
图片
在启动内核函数的时候,需要指定配置参数,它的格式是:
kenerl_function<<<grid_size,block_size>>>
就是告诉CUDA,这次运行的grid的size和block的size,在我们的例子中vectorAdd<<<1,10>>>表示的意思是:Grid中只有一个block,这个block中有10个Thread。
Grid和Block都可以是1维,2维,3维的,这里就不详细描述了。
CUDA生态
前面介绍的是CUDA的冰山一角,希望小伙伴们对CUDA,对GPU编程有个初步认识。
大家也肯定意识到了上面很多cuda开头的各种函数,上层的应用一旦开始使用它们,基本上就和英伟达的CUDA生态绑定了。
图片
在CUDA发展过程中,一个斯坦福的博士生起到了关键作用。
1999年,Nvidia发布了一块叫GeForce的显卡,它的图形处理性能非常出色,非常适合《雷神之锤》游戏。
这时候,斯坦福博士Ian Buck出场了,他疯狂地将32块GeForce显卡连接在一起,再加上8台投影仪,实现了8K分辨率的《雷神之锤》。
玩归玩,他还研究了一下GeForce显卡自带的一个非常原始的编程工具,随后在DARPA的资助下,实现了在GPU上进行通用并行编程。
随后他便加入了英伟达,负责英伟达超级计算包(就是CUDA)的开发。
英伟达的黄教主认为超级计算在未来必将平民化,英伟达要通过CUDA成为领先者。
CUDA的软硬件开发耗资巨大,当2006年正式推出的时候,科技界反应冷淡,认为英伟达瞄准了一个小众的市场,数十亿美元投资有可能打水漂。
英伟达为了推销CUDA,在金融、石油勘探、分子生物等方面孜孜不倦地寻找客户,但都没有起色。
CUDA发展艰难,没有关键应用,缺少重要客户支持。
2008年底,英伟达的股票下跌了70%。
转折点出现在2012年,Hinton团队仅用4个GTX580显卡,利用CUDA技术进行训练出的神经网络,获得了ImageNet比赛的第一名!
机器学习,深度学习彻底被引爆了。
黄仁勋的“赌注”成功了,他在一封邮件中说道:....我们不在是一家GPU公司了,我们是一家AI公司.....
英伟达开始和Google,Facebook等公司合作,推广开源AI框架TensorFlow、PyTorch,当然,它们都构建在CUDA之上。
图片
CUDA彻底统治了AI市场,随后CUDA又发力机器人,自动驾驶等领域。
2023年,以ChatGPT为代表的大模型爆火,英伟达的GPU供不应求,被抢爆了,GPU和CUDA一起攻城掠地,无人可挡。
经过17年的发展,继Windows+Intel , Android + ARM之后,又一个庞大的生态形成了。
这个生态的厉害之处在于:它牢牢占据了软件和硬件的结合之处,CUDA的设计基本就是英伟达硬件形态的抽象。
如果其他GPU厂商想兼容CUDA,就得跟随英伟达的硬件路线,亦步亦趋,相当难受。
如果想重建一套新的生态和API,就会遇到那个老大难问题:软件生态。
英伟达开发了世界上性能最强的GPU,又有着CUDA这个宽广的护城河,照理说,国内厂商是没啥办法的,不用也得用。
但是美国政府送上了神助攻,继A100及H100,连中国专供的“阉割版”A800和H800也不让卖了,禁令甚至波及到了消费级的4090。
原来大家都用英伟达,根本看不上国内产品,现在好了,不得不选国内GPU,比如华为昇腾。
虽然性能差一些,编程接口难用一些,但有总比没有强。
去年11月,百度已经下令将“文心一言”使用的芯片,改向华为芯片,并且为200台服务器购买了1600颗华为昇腾910B AI芯片。
360也表示,采购了华为1,000片左右的AI芯片,和华为合作将AI框架移植到华为昇腾910B的AI芯片。
在实际应用中不断反馈、改善,国产的人工智能芯片肯定会越来越好。
这么发展下去,国内肯定会建立起自己的GPU生态,也会有自己的CUDA。