CUDA內核是如何啓動的？

Question

我已經創建了一個簡單的CUDA應用程序來添加兩個矩陣。它編譯好。我想知道所有線程將如何啓動內核以及CUDA內部的流程如何？我的意思是，每個線程將以什麼方式執行矩陣的每個元素。

我知道這是一個非常基本的概念，但我不知道這一點。我對這個流程感到困惑。

Answer 1

您啓動一個網格塊。

塊被分地分配給多處理器（其中在多處理器塊的數目確定可用共享存儲器的量）。

塊進一步分裂成經紗。對於32個線程的Fermi GPU來說，這些線程執行相同的指令或者是不活動的（因爲它們分支離開，例如通過在同一個warp內退出循環之前的循環，或者沒有使用它們的if）。在Fermi GPU上，最多兩個經線一次在一個多處理器上運行。

每當存在延遲（即執行停止內存訪問或數據依賴關係完成）時，將運行另一個warp（適合一個多處理器的warp數量 - 相同或不同塊 - 由數量決定每個線程使用的寄存器以及一個或多個塊使用的共享內存量）。

此調度發生透明。也就是說，你不必考慮太多。但是，您可能想要使用預定義的整數向量threadIdx（其中是我的線程內的塊？），blockDim（多大一個塊？），blockIdx（其中是我的塊在網格？）和gridDim（多大是網格？）在線程之間拆分工作（讀取：輸入和輸出）。您可能還想了解如何有效地訪問不同類型的內存（因此可以在單個事務中爲多個線程提供服務） - 但這是主題。

NSight提供了圖形化調試器，讓你發生了什麼設備上，一旦你通過行話叢林得到了一個好主意。對於那些在調試器中看不到的東西（例如失速原因或內存壓力），它的分析器也是如此。

您可以通過另一個內核啓動同步網格內的所有線程（全有）。 對於不重疊的順序內核執行，不需要進一步的同步。

一個網格（或一個內核運行 - 無論你怎麼稱呼它）內的線程可以通過使用原子操作（算術）全球內存或適當的存儲器柵欄（用於加載或存儲訪問）進行通信。

您可以使用內部指令__syncthreads()同步所有線程（所有線程之後都會激活 - 儘管一如既往，最多隻能在Fermi GPU上運行兩個線程）。一個塊內的線程可以使用原子操作（用於算術）或適當的內存屏蔽（用於加載或存儲訪問）通過共享或全局內存進行通信。

正如前面提到的，經中的所有線程都始終「同步」，雖然有些可能是無效的。他們可以通過共享或全局存儲器進行通信（或即將推出的具有計算能力的硬件上的「通道交換」3）。您可以使用原子操作（用於算術）和volatile限定的共享或全局變量（加載或存儲訪問在相同的warp內順序進行）。volatile限定符告訴編譯器總是訪問內存，並且永遠不會註冊其他線程無法看到的狀態。

此外，還有全範圍的投票功能，可以幫助您制定分支決策或計算整數（前綴）總和。

好的，基本上就是這樣。希望有所幫助。有一個很好的流動寫作:-)。

Answer 2

嘗試'Cuda-gdb'，這是CUDA調試器。

Answer 3

讓我們除了4×4點矩陣的一個例子..有兩個矩陣A和B，具有尺寸4×4 ..

int main() 
{ 
int *a, *b, *c;   //To store your matrix A & B in RAM. Result will be stored in matrix C 
int *ad, *bd, *cd;   // To store matrices into GPU's RAM. 
int N =4;     //No of rows and columns. 

size_t size=sizeof(float)* N * N; 

a=(float*)malloc(size);  //Allocate space of RAM for matrix A 
b=(float*)malloc(size);  //Allocate space of RAM for matrix B 

//allocate memory on device 
    cudaMalloc(&ad,size); 
    cudaMalloc(&bd,size); 
    cudaMalloc(&cd,size); 

//initialize host memory with its own indices 
    for(i=0;i<N;i++) 
     { 
    for(j=0;j<N;j++) 
     { 
      a[i * N + j]=(float)(i * N + j); 
      b[i * N + j]= -(float)(i * N + j); 
     } 
     } 

//copy data from host memory to device memory 
    cudaMemcpy(ad, a, size, cudaMemcpyHostToDevice); 
    cudaMemcpy(bd, b, size, cudaMemcpyHostToDevice); 

//calculate execution configuration 
    dim3 grid (1, 1, 1); 
    dim3 block (16, 1, 1); 

//each block contains N * N threads, each thread calculates 1 data element 

    add_matrices<<<grid, block>>>(ad, bd, cd, N); 

    cudaMemcpy(c,cd,size,cudaMemcpyDeviceToHost); 
    printf("Matrix A was---\n"); 
    for(i=0;i<N;i++) 
    { 
     for(j=0;j<N;j++) 
      printf("%f ",a[i*N+j]); 
     printf("\n"); 
    } 

    printf("\nMatrix B was---\n"); 
    for(i=0;i<N;i++) 
    { 
     for(j=0;j<N;j++) 
      printf("%f ",b[i*N+j]); 
     printf("\n"); 
    } 

    printf("\nAddition of A and B gives C----\n"); 
    for(i=0;i<N;i++) 
    { 
     for(j=0;j<N;j++) 
      printf("%f ",c[i*N+j]); //if correctly evaluated, all values will be 0 
     printf("\n"); 
    } 



    //deallocate host and device memories 
    cudaFree(ad); 
    cudaFree(bd); 
    cudaFree (cd); 

    free(a); 
    free(b); 
    free(c); 

    getch(); 
    return 1; 
} 

/////Kernel Part 

__global__ void add_matrices(float *ad,float *bd,float *cd,int N) 
{ 
    int index; 
    index = blockIDx.x * blockDim.x + threadIDx.x    

    cd[index] = ad[index] + bd[index]; 
} 

讓我們加入16×16點的矩陣的一個例子.. 你有兩個矩陣A和B，尺寸爲16 * 16 ..

首先你必須決定你的線程配置。 假設您將啓動一個內核函數，它將執行矩陣加法的並行計算，這將在您的GPU設備上執行。

現在，一個網格啓動時只有一個內核函數。 一個網格最多可以有65,535個塊，它們可以以三維方式排列。 （65535 * 65535 * 65535）。

在網格中的每個塊可以有最高1024不threads.Those線程也可以安排在3點維的方式（1024 * 1024 * 64）

現在我們的問題是，除了16×16點的矩陣..

A | 1 2 3 4 |  B | 1 2 3 4 |  C| 1 2 3 4 | 
    | 5 6 7 8 | +  | 5 6 7 8 | = | 5 6 7 8 | 
    | 9 10 11 12 |   | 9 10 11 12 |  | 9 10 11 12 | 
    | 13 14 15 16|   | 13 14 15 16|  | 13 14 15 16| 

我們需要16個線程來執行計算。

i.e. A(1,1) + B (1,1) = C(1,1) 
    A(1,2) + B (1,2) = C(1,2) 
    .  .   . 
    .  .   . 
    A(4,4) + B (4,4) = C(4,4) 

所有這些線程將同時執行。 所以我們需要一個有16個線程的塊。 爲了方便起見，我們將在一個塊中安排線程（16 * 1 * 1） 由於線程數不是16，所以我們只需要一個塊來存儲這16個線程。

這樣，網格配置將是dim3 Grid(1,1,1)即網格將僅具有一個塊 和嵌段構型將是dim3 block(16,1,1)即塊將具有佈置逐列16個線程。

以下程序會給你清晰的執行概念。 瞭解索引部分（即threadIDs，blockDim，blockID）是重要的部分。您需要閱讀CUDA文獻。一旦你對索引有了清晰的認識，你就會贏得半場戰鬥！因此，花一些時間與cuda書籍，不同的算法和紙筆當然！