初識cuda一文通

寒柏懒得想發表於2024-11-03

cuda學習部落格

本文為本人cuda學習過程中的記錄和理解，多參考@譚升等大佬前輩的部落格，以及NVIDIA官方文件。如有錯誤煩請指正，如有侵權請聯絡刪除。

0. 平行計算與計算機架構

計算機架構是平行計算或者說HPC領域中十分重要的內容。

。。看書後期記得補充。。

無論是OpenMP還是CUDA，寫並行程式主要是分解任務，在軟體層面上大體可以分為‘指令並行’和‘資料並行’。

1. 異構計算與CUDA

cuda基礎

進入正題，我們從hello world開始學習cuda

#include<stdio.h>
__global__ void hello_world(void)
{
    printf("hello world!);
}

int main(int argc, char **argv)
{
    hello_world<<<1, 10>>>();
    cudaDeviceReset();
    return 0;
}

下面我們具體審視一下這個簡單的cuda程式。
（1）__global__字首，cuda的精髓，即核函式。核函式的程式碼將在成千上萬個執行緒上執行。
（2）cudaDeviceReset(),實現CPU和Device（即GPU）的同步，保證cpu和gpu一起退出程式。
大家這裡先不用著急，後續我們還會具體談cuda中執行單元的劃分，並NVIDIA為我們提供的若干工具來分析編寫並行程式。

在異構環境中，主機host和裝置device透過pcie匯流排通訊。通訊內容就包括了最重要的記憶體資訊傳遞。CUDA上的記憶體管理API大體和cpu中相同，例如：
cudaMalloc()諸如此類

跑在匯流排上的一個重要API：cudaMemcpy()記憶體複製函式，可以實現DeviceToHost/HostTODevice等等過程。

GPU架構

執行緒管理

下面我們進入到執行緒管理。首先，一個核函式只能有一個grid，每個grid可以分為多個block，每個block又可以分為多個thread。CUDA還為同一個block中的執行緒提供了共享記憶體機制，在每個block中單獨劃分出一塊記憶體空間用於共享，透過關鍵詞__shared__訪問，擁有比全域性記憶體更好的效能表現，這點後續我們也會詳談。

核函式

kernel_name<<<grid,block>>>
注意host端核函式執行是非同步的，並且當主機發出執行核函式的指令後，會立即收回控制權。此刻就需要方法來使host和device同步，最常見的是使用cudaMemcpy來使host等待device資料，從而實現同步。

在CUDA程式編寫時，最常見的是將序列程式中的for迴圈並行化（這點在OpenMP和MPI中已經很熟悉了）

細談執行緒束Warp

CUDA中的同步

無論是pthread中的🔓還是openmp，都有相應的同步機制，CUDA為避免記憶體競爭，也有同步方法。__syncthread()可以實現同block內的執行緒同步，想要同步不同block的執行緒，方法是藉助核函式執行。

CUDA並行效能分析

這裡我們使用簡單的矩陣加法為例來進行實驗分析，硬體平臺為RTX3060

分析工具介紹

  • nvprof：

動態並行

動態並行或者叫他巢狀並行，類似於父程序/子程序的概念。

2.CUDA記憶體

  • 暫存器
  • 共享記憶體
  • 常量記憶體
  • 全域性記憶體
  • 紋理記憶體
  • 本地記憶體

記憶體管理

cudaMelloc() //分配記憶體
cudaMemset() //初始化記憶體
cudaFree() //釋放記憶體
cudaemcpy() //傳輸記憶體

特殊記憶體

固定記憶體：固定記憶體傳輸速更快，但是分配和釋放成本更高

零複製記憶體：
統一記憶體定址：cudaMallocManaged()

記憶體訪問模式

這部分內容至關重要，對於程式調優來說，記得程式碼驗證

1.對齊合併記憶體訪問，減少頻寬浪費
2.足夠的併發記憶體操作，隱藏記憶體延遲

共享記憶體

相關文章