厲害,美國人又搞了一個壟斷的生態系統
去年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。