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

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

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

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

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

CPU vs GPU

圖片

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

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

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

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

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

(1) ALU 簡單，但是超級多

(2) Cache很小

(3) 邏輯控制簡單。

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

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

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

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

把串行改成并行

我們用一個例子來展示一下：

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在執(zhí)行時會做一個循環(huán)，然后把兩個數(shù)組對應(yīng)的元素進(jìn)行相加，結(jié)果存到數(shù)組c中。

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

如何把它改成并行計算呢？

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

圖片

程序員該怎么寫代碼，來表達(dá)這個想法呢？

CUDA

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

圖片

在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分別是兩個要想加的數(shù)組，c用來保存結(jié)果。

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

threadIdx是個什么東西？

似乎是個線程的索引，找到這個線程的index以后，取出a,b中index對應(yīng)的值，加起來放到c中。例如index是0，那就取出a[0],b[0]加起來，放到c[0]中，這就實現(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個參數(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ù)，這里指定了一組10個線程(數(shù)組長度為10)。

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

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

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

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

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

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

圖片

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

圖片

多個Block又形成一個Grid，被映射到整個CPU

圖片

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

kenerl_function<<<grid_size,block_size>>>

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

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

CUDA生態(tài)

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

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

圖片

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

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

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

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

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

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

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

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

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

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

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

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

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

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

圖片

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

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

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

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

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

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

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

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

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

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

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

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

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

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