CUDA在執行期間結合了線程獨立（??）變量

如果標題混亂，我很抱歉。我雖然漫長而艱苦，並且無法想出合適的方式來將問題單獨列出來。所以這裏有更多的細節。我正在做一個基本的圖像減法，其中第二個圖像已被修改，我需要找出多少變化與圖像的比例。爲此，我使用了下面的代碼。這兩個圖像都是128x1024。CUDA在執行期間結合了線程獨立（??）變量

for(int i = 0; i < 128; i++) 
{ 
    for(int j = 0; j < 1024; j++) 
    { 
     den++; 
     diff[i * 1024 + j] = orig[i * 1024 + j] - modified[i * 1024 + j]; 
     if(diff[i * 1024 + j] < error) 
     { 
      num++; 
     } 
    } 
} 
ratio = num/den;

上面的代碼工作在CPU上不錯，但我想嘗試這樣做的CUDA。爲此，我可以設置CUDA執行圖像的基本減法（下面的代碼），但我無法弄清楚如何執行條件if語句來獲取我的比例。

__global__ void calcRatio(float *orig, float *modified, int size, float *result) 
{ 
    int index = threadIdx.x + blockIdx.x * blockDim.x; 
    if(index < size) 
     result[index] = orig[index] - modified[index]; 
}

所以，到現在爲止它的工作原理，但我無法弄清楚如何parrallelize的Num和書房計數器在每個線程的所有線程執行結束來計算比例。對我來說，感覺就像num和den counders是獨立於線程的，因爲每次我嘗試使用它們時，似乎它們只會增加一次。

任何幫助將不勝感激，因爲我剛剛開始使用CUDA，並且我在網上看到的每個例子似乎都不適用於我需要做的事情。

編輯：修復了我的天真代碼。忘了輸入代碼中的主要條件之一。這是漫長的一天。

for(int i = 0; i < 128; i++) 
{ 
    for(int j = 0; j < 1024; j++) 
    { 
     if(modified[i * 1024 + j] < 400.0) //400.0 threshold value to ignore noise 
     { 
      den++; 
      diff[i * 1024 + j] = orig[i * 1024 + j] - modified[i * 1024 + j]; 
      if(diff[i * 1024 + j] < error) 
      { 
       num++; 
      } 
     } 
    } 
} 
ratio = num/den;

來源

2011-04-14 user656210

您需要使用跨所有線程執行全局總和的操作稱爲「並行減少」。雖然你可以使用原子操作來做到這一點，但我不會推薦它。有一個減少內核和一個非常好的論文討論CUDA SDK中的技術，值得一讀。

如果我寫代碼，做你想做的，它可能是這樣的：

template <int blocksize> 
__global__ void calcRatio(float *orig, float *modified, int size, float *result, 
          int *count, const float error) 
{ 
    __shared__ volatile float buff[blocksize]; 

    int index = threadIdx.x + blockIdx.x * blockDim.x; 
    int stride = blockDim.x * gridDim.x; 

    int count = 0; 
    for(int i=index; i<n; i+=stride) { 
     val = orig[index] - modified[index]; 
     count += (val < error); 
     result[index] = val; 
    } 

    buff[threadIdx.x] = count; 
    __syncthreads(); 


    // Parallel reduction in shared memory using 1 warp 
    if (threadId.x < warpSize) { 

     for(int i=threadIdx.x + warpSize; i<blocksize; i+= warpSize) { 
      buff[threadIdx.x] += buff[i]; 

     if (threadIdx.x < 16) buff[threadIdx.x] +=buff[threadIdx.x + 16]; 
     if (threadIdx.x < 8) buff[threadIdx.x] +=buff[threadIdx.x + 8]; 
     if (threadIdx.x < 4) buff[threadIdx.x] +=buff[threadIdx.x + 4]; 
     if (threadIdx.x < 2) buff[threadIdx.x] +=buff[threadIdx.x + 2]; 
     if (threadIdx.x == 0) count[blockIdx.x] = buff[0] + buff[1]; 
    } 
}

第一節中做您的串行代碼做什麼 - 計算的差異和線程局部共小於錯誤的元素。注意我已經編寫了這個版本，以便每個線程都可以處理多個輸入數據。這樣做的目的是幫助抵消後續並行減少的計算成本，並且想法是使用的塊和線程數少於輸入數據集條目的數量。

第二節是本身的減少，在共享內存中完成。它實際上是一個「類似樹」的操作，其中一個線程塊內的線程局部小計集合的大小首先被總計爲32個小計，然後合併小計直到存在該塊的最終小計，並且然後被存儲的總數爲塊。您將得到一個小計的小計清單，您可以將每個啓動的塊複製一份，然後將其複製回主機，並計算出您需要的最終結果。

請注意我在瀏覽器中對此進行了編碼，未編譯它，可能存在錯誤，但它應該提供一個有關您正在嘗試執行的「高級」版本如何工作的想法。

來源

2011-04-14 05:43:43 talonmies

感謝talonmies。但我只是意識到我在打字代碼中犯了一個巨大的錯誤。我在記憶中漫長的一天之後輸入了這個內容，因此忘記了一個主要部分。上面編輯我的帖子。 – user656210 2011-04-14 18:20:46

這對代碼幾乎沒有影響。您添加的條件在我發佈的內核的第一個節的循環內進行。 – talonmies 2011-04-14 18:56:06

Ooooo。這大概是我在想什麼......但表達得更清楚。你知道這個操作的名字。這可能在未來有用。 – cgmb 2011-04-15 03:27:29

分母非常簡單，因爲它只是大小。

分子比較麻煩，因爲給定線程的值取決於以前的所有值。你將不得不連續做這個操作。

你正在尋找的東西可能是atomicAdd。雖然這很慢。

我想你會發現這個問題有關。你的數字基本上是全球數據。 CUDA array-to-array sum

或者，您可以將錯誤檢查的結果轉儲到數組中。計算結果可以並行化。這會有點棘手，但我認爲這樣會擴大：http://tekpool.wordpress.com/2006/09/25/bit-count-parallel-counting-mit-hakmem/

來源

2011-04-14 01:18:23 cgmb

CUDA在執行期間結合了線程獨立（??）變量

回答

相關問題