去年AIGC大火，程序員都把注意力放在了最上層，而忽略了提供算力的最底層：GPU。

不過(guò)這也正常，就像很少人直接針對(duì)CPU編程一樣，直接針對(duì)GPU編程的人也不多。

但是了解一下GPU編程，絕對(duì)大有好處。

今天先聊聊GPU編程，然后再聊聊一個(gè)CUDA這個(gè)新的生態(tài)系統(tǒng)，對(duì)編程細(xì)節(jié)不感興趣的可以直接拉到最后。

對(duì)了，文末還有免費(fèi)送書(shū)的福利。

CPU vs GPU

圖片

CPU的設(shè)計(jì)目標(biāo)是“盡可能地降低延時(shí)”

(1) 強(qiáng)大的ALU（算術(shù)邏輯單元），可以在很少的時(shí)鐘周期內(nèi)完成算術(shù)運(yùn)算。

(2) 巨大的Cache：加快指令和數(shù)據(jù)的存取速度

(3) 復(fù)雜的邏輯控制：當(dāng)程序員有多個(gè)分支，它可以通過(guò)分支預(yù)測(cè)來(lái)降低延時(shí)。

GPU的目標(biāo)是：“盡可能地實(shí)現(xiàn)大吞吐量”

(1) ALU 簡(jiǎn)單，但是超級(jí)多

(2) Cache很小

(3) 邏輯控制簡(jiǎn)單。

如果把GPU的單個(gè)核心比作小學(xué)生，那一個(gè)CPU的核心就是老教授。

如果要做微積分，幾千個(gè)小學(xué)生也比如上老教授。

但是，如果只是100以?xún)?nèi)的加減法，幾千個(gè)小學(xué)生同時(shí)做（并行計(jì)算），那效率肯定要比老教授高。

老教授處理復(fù)雜任務(wù)的能力是碾壓小學(xué)生的，但是對(duì)于沒(méi)有那么復(fù)雜的任務(wù)，還是頂不住人多。

把串行改成并行

我們用一個(gè)例子來(lái)展示一下：

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;
}

這段簡(jiǎn)單的代碼大家都能看懂，CPU在執(zhí)行時(shí)會(huì)做一個(gè)循環(huán)，然后把兩個(gè)數(shù)組對(duì)應(yīng)的元素進(jìn)行相加，結(jié)果存到數(shù)組c中。

由于是順序處理的，如果數(shù)組非常大，就會(huì)比較耗時(shí)。

如何把它改成并行計(jì)算呢？

數(shù)組中有10個(gè)元素，我們可以創(chuàng)建10個(gè)線(xiàn)程，把每個(gè)線(xiàn)程扔到一個(gè)GPU核心中去運(yùn)行。

圖片

程序員該怎么寫(xiě)代碼，來(lái)表達(dá)這個(gè)想法呢？

CUDA

英偉達(dá)的CUDA是一個(gè)并行計(jì)算平臺(tái)，可以讓程序員可以通過(guò)C、C++等語(yǔ)言在GPU上并行執(zhí)行代碼。

圖片

在CUDA中，把CPU所在的部分叫做Host，GPU稱(chēng)為Device，它們之間通過(guò)總線(xiàn)相連。

圖片

對(duì)于之前的例子，CUDA代碼是這樣的：

__global__ void vectorAdd(int* a, int* b, int* c){
    int i = threadIdx.x;
    c[i] = a[i] + b[i];
    return;
}

估計(jì)大部分小伙伴都能猜出來(lái)這段代碼的含義。

a,b分別是兩個(gè)要想加的數(shù)組，c用來(lái)保存結(jié)果。

__global__應(yīng)該是個(gè)指示符，表示這段代碼是個(gè)“內(nèi)核函數(shù)”，要被放到GPU上來(lái)執(zhí)行。

threadIdx是個(gè)什么東西？

似乎是個(gè)線(xiàn)程的索引，找到這個(gè)線(xiàn)程的index以后，取出a,b中index對(duì)應(yīng)的值，加起來(lái)放到c中。例如index是0，那就取出a[0],b[0]加起來(lái)，放到c[0]中，這就實(shí)現(xiàn)了我們之前的想法。

值得注意的是，這里的a,b,c不是Host的內(nèi)存，而是Device（GPU）的內(nèi)存，所以我們得把原始的數(shù)據(jù)復(fù)制到GPU中。

1. 先在GPU中分配內(nèi)存

int* cudaA = 0;
int* cudaB = 0;
int* cudaC = 0;
// 使用cudaMalloc在GPU中分配內(nèi)存
cudaMalloc(&cudaA,sizeof(a));
cudaMalloc(&cudaB,sizeof(b));
cudaMalloc(&cudaC,sizeof(c));

2.然后把原始數(shù)據(jù)從Host復(fù)制到Device(即GPU)中

//注意第4個(gè)參數(shù)，是從Host 到 Device
cudaMemcpy(cudaA, a, sizeof(a), cudaMemcpyHostToDevice);
cudaMemcpy(cudaB, b, sizeof(b), cudaMemcpyHostToDevice);

3. 調(diào)用內(nèi)核函數(shù)

vectorAdd <<<1, sizeof(a) / sizeof(a[0])>>> (cudaA, cudaB, cudaC);

調(diào)用vectorAdd的時(shí)候，被<<< >>>包圍起來(lái)的部分是配置參數(shù)，這里指定了一組10個(gè)線(xiàn)程(數(shù)組長(zhǎng)度為10)。

這10個(gè)線(xiàn)程會(huì)被放到10個(gè)GPU核心中去執(zhí)行，他們的索引是從0到9。

所以在vectorAdd函數(shù)中可以通過(guò)threadIdx.x引用到當(dāng)前線(xiàn)程的索引，例如9 ， 那就知道當(dāng)前線(xiàn)程要做的事情：把a(bǔ)[9]和b[9]加起來(lái)，放到c[9]中。

這樣10個(gè)GPU核心就是同時(shí)執(zhí)行10次加法，速度飛快。

4. 把結(jié)果復(fù)制回Host

// 注意第4個(gè)參數(shù)，是從Device 到 Host
cudaMemcpy(c, cudaC, sizeof(c), cudaMemcpyDeviceToHost);

小伙伴們肯定已經(jīng)意識(shí)到了，這里邊有個(gè)核心的概念：Thread（線(xiàn)程），每個(gè)線(xiàn)程都會(huì)被映射到一個(gè)GPU核心去執(zhí)行。

圖片

多個(gè)Thread可以組成一個(gè)塊（Block），被映射到多個(gè)核心

圖片

多個(gè)Block又形成一個(gè)Grid，被映射到整個(gè)CPU

圖片

在啟動(dòng)內(nèi)核函數(shù)的時(shí)候，需要指定配置參數(shù)，它的格式是：

kenerl_function<<<grid_size,block_size>>>

就是告訴CUDA，這次運(yùn)行的grid的size和block的size，在我們的例子中vectorAdd<<<1,10>>>表示的意思是：Grid中只有一個(gè)block，這個(gè)block中有10個(gè)Thread。

Grid和Block都可以是1維，2維，3維的，這里就不詳細(xì)描述了。

CUDA生態(tài)

前面介紹的是CUDA的冰山一角，希望小伙伴們對(duì)CUDA，對(duì)GPU編程有個(gè)初步認(rèn)識(shí)。

大家也肯定意識(shí)到了上面很多cuda開(kāi)頭的各種函數(shù)，上層的應(yīng)用一旦開(kāi)始使用它們，基本上就和英偉達(dá)的CUDA生態(tài)綁定了。

圖片

在CUDA發(fā)展過(guò)程中，一個(gè)斯坦福的博士生起到了關(guān)鍵作用。

1999年，Nvidia發(fā)布了一塊叫GeForce的顯卡，它的圖形處理性能非常出色，非常適合《雷神之錘》游戲。

這時(shí)候，斯坦福博士Ian Buck出場(chǎng)了，他瘋狂地將32塊GeForce顯卡連接在一起，再加上8臺(tái)投影儀，實(shí)現(xiàn)了8K分辨率的《雷神之錘》。

玩歸玩，他還研究了一下GeForce顯卡自帶的一個(gè)非常原始的編程工具，隨后在DARPA的資助下，實(shí)現(xiàn)了在GPU上進(jìn)行通用并行編程。

隨后他便加入了英偉達(dá)，負(fù)責(zé)英偉達(dá)超級(jí)計(jì)算包（就是CUDA）的開(kāi)發(fā)。

英偉達(dá)的黃教主認(rèn)為超級(jí)計(jì)算在未來(lái)必將平民化，英偉達(dá)要通過(guò)CUDA成為領(lǐng)先者。

CUDA的軟硬件開(kāi)發(fā)耗資巨大，當(dāng)2006年正式推出的時(shí)候，科技界反應(yīng)冷淡，認(rèn)為英偉達(dá)瞄準(zhǔn)了一個(gè)小眾的市場(chǎng)，數(shù)十億美元投資有可能打水漂。

英偉達(dá)為了推銷(xiāo)CUDA，在金融、石油勘探、分子生物等方面孜孜不倦地尋找客戶(hù)，但都沒(méi)有起色。

CUDA發(fā)展艱難，沒(méi)有關(guān)鍵應(yīng)用，缺少重要客戶(hù)支持。

2008年底，英偉達(dá)的股票下跌了70%。

轉(zhuǎn)折點(diǎn)出現(xiàn)在2012年，Hinton團(tuán)隊(duì)僅用4個(gè)GTX580顯卡，利用CUDA技術(shù)進(jìn)行訓(xùn)練出的神經(jīng)網(wǎng)絡(luò)，獲得了ImageNet比賽的第一名！

機(jī)器學(xué)習(xí)，深度學(xué)習(xí)徹底被引爆了。

黃仁勛的“賭注”成功了，他在一封郵件中說(shuō)道：....我們不在是一家GPU公司了，我們是一家AI公司.....

英偉達(dá)開(kāi)始和Google,Facebook等公司合作，推廣開(kāi)源AI框架TensorFlow、PyTorch，當(dāng)然，它們都構(gòu)建在CUDA之上。

圖片

CUDA徹底統(tǒng)治了AI市場(chǎng)，隨后CUDA又發(fā)力機(jī)器人，自動(dòng)駕駛等領(lǐng)域。

2023年，以ChatGPT為代表的大模型爆火，英偉達(dá)的GPU供不應(yīng)求，被搶爆了，GPU和CUDA一起攻城掠地，無(wú)人可擋。

經(jīng)過(guò)17年的發(fā)展，繼Windows+Intel , Android + ARM之后，又一個(gè)龐大的生態(tài)形成了。

這個(gè)生態(tài)的厲害之處在于：它牢牢占據(jù)了軟件和硬件的結(jié)合之處，CUDA的設(shè)計(jì)基本就是英偉達(dá)硬件形態(tài)的抽象。

如果其他GPU廠(chǎng)商想兼容CUDA，就得跟隨英偉達(dá)的硬件路線(xiàn)，亦步亦趨，相當(dāng)難受。

如果想重建一套新的生態(tài)和API，就會(huì)遇到那個(gè)老大難問(wèn)題：軟件生態(tài)。

英偉達(dá)開(kāi)發(fā)了世界上性能最強(qiáng)的GPU，又有著CUDA這個(gè)寬廣的護(hù)城河，照理說(shuō)，國(guó)內(nèi)廠(chǎng)商是沒(méi)啥辦法的，不用也得用。

但是美國(guó)政府送上了神助攻，繼A100及H100，連中國(guó)專(zhuān)供的“閹割版”A800和H800也不讓賣(mài)了，禁令甚至波及到了消費(fèi)級(jí)的4090。

原來(lái)大家都用英偉達(dá)，根本看不上國(guó)內(nèi)產(chǎn)品，現(xiàn)在好了，不得不選國(guó)內(nèi)GPU，比如華為昇騰。

雖然性能差一些，編程接口難用一些，但有總比沒(méi)有強(qiáng)。

去年11月，百度已經(jīng)下令將“文心一言”使用的芯片，改向華為芯片，并且為200臺(tái)服務(wù)器購(gòu)買(mǎi)了1600顆華為昇騰910B AI芯片。

360也表示，采購(gòu)了華為1,000片左右的AI芯片，和華為合作將AI框架移植到華為昇騰910B的AI芯片。

在實(shí)際應(yīng)用中不斷反饋、改善，國(guó)產(chǎn)的人工智能芯片肯定會(huì)越來(lái)越好。

這么發(fā)展下去，國(guó)內(nèi)肯定會(huì)建立起自己的GPU生態(tài)，也會(huì)有自己的CUDA。