2014-02-12 60 views
3

我對線程的形成和執行方式有很多疑問。CUDA線程如何工作

首先,文檔將GPU線程描述爲輕量級線程。 假設我希望乘以兩個100*100矩陣。如果每個元素都由不同的線程計算,這將需要100*100個線程。但是,我的GPU(NVIDIA GT 640M LE)規格顯示兩個SM,每個SM只能支持2048個線程。如果我的GPU不支持這麼多的線程,那麼如何計算其餘元素並行呢?

另外考慮基本向量添加代碼。假設我調用與1塊和64個線程內核添加100個元件的兩個陣列,每個陣列如下:

__global__ void add(int* a,int* b,int* c) 
    { 
     int i = threadIdx.x; 
     for(i<100) 
     { 
      c[i] = a[i] + b[i]; 
     {  
    } 

由於只有64個線程進行初始化我假定64個元件被並行加入。

  • 如何添加其餘元素?
  • warp調度程序如何決定分配哪些線程來添加最後的36個元素?

我的主要問題是:

我不明白一個線程是如何知道哪些元素進行操作。

回答

6

您的卡具有計算能力3.0,請參閱here

從CUDA C編程指南的表12中可以看出,您爲計算能力提及的2048線程的數量是指每個多處理器的最大駐留線程數。這並不意味着您無法在整體上啓動超過2048的線程。例如,從該表的上方几行可以看出,線程塊網格的最大最大值爲2^31-1。這意味着,例如,啓動8192線程的1d線程網格是完全合法的。原因是該卡將執行線程扭曲之間的上下文切換,如本文所示:What is the context switching mechanism in GPU?

關於你的問題的第二部分,你實現add函數在概念上是錯誤的。您正在使用索引i作爲線索索引和作爲for循環索引。更正確的實現是以下

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

上述寫入裝置執行以下操作:每個線程將執行兩個任務,即現在

int i = threadIdx.x; 
    c[i] = a[i] + b[i]; 

,例如,用於螺紋#3threadIdx.x的值變量將是3。因此,線程#3將處理一個局部變量i,它的內存空間是私有的,其值將被分配給3。此外,它將從全局內存中加載a[3]b[3],將它們加起來,將結果分配給c[3],然後將最終結果存儲到全局內存。因此,當您啓動網格時,您當然不能僅用64線程填充整個100元素的數組,並且您將需要100線程。

請注意,以上解釋過於簡單。我建議你閱讀一些基本的教科書作爲着名的CUDA示例。

+0

感謝您的答案。我瞭解有關啓動大量線程的部分。 – mastercheif141

+0

即將到來的第二部分我還沒有得到如何添加其餘的元素。因爲只有64個線程被啓動,所以在該塊上僅存在兩個warp。在添加了64個元素之後,將添加剩餘的元素?如果是,那麼通過哪些線程? – mastercheif141

+0

@ mastercheif141我試圖更好地解釋你的第二個問題。見編輯的答案。 – JackOLantern

1

會給你CUDA中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 =16; 

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

讓我們有兩個矩陣A和B添加16×16點的矩陣.. 的一個例子,其尺寸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書籍:-)

0

因爲是非常錯誤在這裏 - 一些線程與螺紋編號< 100將運行forewer。 對於新手來說,可以這樣解釋:threadid是由系統值預先定義的,顯示當前線程號碼。當前線程需要它從一個值,從B,用C寫它,所以它會

int i = threadIdx.x; 
c[i] = a[i] + b[i]; 

如果你有數組大小魔神100什麼不是爲了某個線程匹配的64倍, 塊大小不讀/寫超出範圍,請執行以下操作:

int i = threadIdx.x; 
    if(i < 100){ 

     c[i] = a[i] + b[i]; 
    } 

只有在最後一個塊上纔有分歧。 也許你想