如何計算塊編號

我寫一個CUDA代碼，我使用蘋果牛9500 GT顯卡。如何計算塊編號

我試圖處理20000000個整數元素的陣列，並且我使用的線程數爲256

warp大小爲32。計算能力爲1.1

這是硬件http://www.geforce.com/hardware/desktop-gpus/geforce-9500-gt/specifications

現在，block num = 20000000/256 = 78125？

這種聲音不正確。我如何計算塊號？任何幫助，將不勝感激。

我的CUDA內核函數如下。這個想法是每個塊會計算它的總和，然後最後的總和將通過加上每個塊的總和來計算。

__global__ static void calculateSum(int * num, int * result, int DATA_SIZE) 
{ 
    extern __shared__ int shared[]; 
    const int tid = threadIdx.x; 
    const int bid = blockIdx.x; 

    shared[tid] = 0; 
    for (int i = bid * THREAD_NUM + tid; i < DATA_SIZE; i += BLOCK_NUM * THREAD_NUM) { 
     shared[tid] += num[i]; 
    } 

    __syncthreads(); 
    int offset = THREAD_NUM/2; 
    while (offset > 0) { 
     if (tid < offset) { 
      shared[tid] += shared[tid + offset]; 
     } 
     offset >>= 1; 
     __syncthreads(); 
    } 

    if (tid == 0) { 
     result[bid] = shared[0]; 

    } 
}

而且我把這種功能

calculateSum <<<BLOCK_NUM, THREAD_NUM, THREAD_NUM * sizeof(int)>>> (gpuarray, result, size);

其中THREAD_NUM = 256 和GPU數組大小的2000萬

在這裏，我只是用塊編號爲16，但不知道這是正確的？如何確保達到最大並行度？

這裏是我的CUDA佔用計算器的輸出。它說，當塊號爲8時，我將擁有100％的佔用率。這意味着當塊號= 8和線程號= 256時，我將獲得最大效率。那是對的嗎？

CUDA Occupancy calculation 感謝

來源

2012-05-02 Coder

您錯誤地解釋了佔用計算器的輸出。它表示每個多處理器的最佳塊數爲3（第18行）。因此（在這種情況下），每個多處理器* 4多處理器= 12塊需要3個塊才能實現該內核的最優並行性*。 – talonmies

如果每個thred過程中的一個元素，每塊有256個線程，你應該運行20000000個線程，從而準確地78125塊。這是完全有效的數字。

但是，有一點問題。我沒有CC1.1設備在手，但在CC1.3：

Maximum sizes of each dimension of a grid:  65535 x 65535 x 1

所以您應該運行的內核多次爲數據的不同部分，或使2D網格，只是平凡變換分析2D地址線程到數組元素的1D地址。

來源

2012-05-02 12:17:26 aland

感謝您的輸入我剛纔發佈了我的內核代碼。你能否請檢查它的正確 – Coder

我認爲對於CC 1.1最大尺寸是65535 * 65535 * 1根據這個 http://en.wikipedia.org/wiki/CUDA#Version_features_and_specifications – Coder

在你的情況下，線程總數（20000000）均勻地除以每個塊的線程數（256），所以你可以使用該數字（78125）。如果這些數字不能均勻分配，則規則的整數除法會將其舍入，最終會得到比您需要的線程更少的線程。因此，在這種情況下，你需要四捨五入的分歧的結果了類似以下的功能：

int DivUp(int a, int b) { 
    return ((a % b) != 0) ? (a/b + 1) : (a/b); 
}

由於此功能可能會給你更多的線程之外還有元素，則還需要添加一個測試在你的內核中放棄最後幾個線程的計算：

int i(blockIdx.x * blockDim.x + threadIdx.x); 
if (i >= n_items) { 
    return; 
}

但是，還有一個額外的障礙。您的硬件在網格中的每個維度上限制爲最多65535個塊，並且僅限於兩個維度（x和y）。因此，如果在使用DivUp（）之後，最終的計數值會高於此值，那麼您有兩種選擇。您可以分割工作負載並多次運行內核，也可以使用兩個維度。

要使用兩個維度，請選擇兩個數字，每個數字都低於硬件限制，並且在乘以時會成爲您需要的實際塊數。然後，在內核頂部添加代碼，將兩個維度（x和y）合併到一個索引中。

來源

2012-05-02 12:35:44

您可以請評論我的代碼上面。謝謝 – Coder

您發佈的內核代碼可以處理任何輸入數據大小，與您選擇啓動的塊數無關。選擇應該簡單到性能。

作爲一個經驗法則，對於這種類型的內核，您希望儘可能多的塊在單個多處理器上同時運行，使卡上的多處理器數量增加一倍。第一個數字可以使用CUDA工具包中的CUDA佔用電子表格獲得，但每個多處理器的上限爲8個塊，第二個數字爲的設備爲4 。這意味着不需要超過32個塊來實現最大可能的並行性，但要準確回答需要訪問我目前沒有的編譯器。

您還可以使用基準測試通過4,8,12,16,20,24,28或32個塊之一（4的倍數）以實驗方式確定最佳塊數，因爲這是您的多處理器的數量卡）。

來源

2012-05-02 17:50:35 talonmies

我真的看到了你的答案，並且當我再次與Cuda一起工作時，我將定義測試它。如果使用更少的塊，並且因此每個線程中處理的元素更多，速度更快，爲什麼總是存在3D網格和塊結構，並且幾乎每本書和源都說明儘可能多地使用線程，因爲它是SIMD體系結構。來自其原始着色歷史的壞習慣？ – djmj

感謝您的輸入。我附上了我的cuda佔用率計算器輸出。我的理解是否正確？ – Coder

由於您擁有3個駐留塊，256個線程等於每個流式多處理器（SM）的最大768個駐留線程，因此您將自動擁有100％的佔用率。這是每個SM的並行處理線程數。佔用率由每個SM的駐留塊和線程定義。 Aslong所有的SM都很忙，你將有100％的入住率。如果你的線程總數是768（例子96），你總是擁有100％的佔用率。 – djmj

您只在內核中使用網格的x維。所以你僅限於使用cc 1.1的65535個塊。

20000000/256 = 78125是正確的！

所以你肯定需要多一個塊。

內核：

//get unique block index 
const unsigned int blockId = blockIdx.x //1D 
    + blockIdx.y * gridDim.x //2D 

//terminate unnecessary blocks 
if(blockId >= 78124) 
    return; 

//... rest of kernel

最簡單的方法是使用兩個Y塊和校驗塊ID在內核。

dim3 gridDim = dim3(65535, 2);

這將使更多的則52945塊也沒用，我不知道什麼的開銷，但填補了第一個X，然後y和z維度可以創造很多未使用的塊，特別是如果達到z方向！

（Nvidia的應definetly已經提供了效用函數，得到內部內核獨特塊使用的最佳網格的使用作爲其在這裏的情況下）

對於這個簡單的例子，如何有關使用x和y，並計算根。

grid(280, 280) = 78400 blocks //only 275 blocks overhead, less is not possible

這是計算能力3.0的一大優勢。每個塊上的32位範圍使得生活通常更容易。爲什麼它被限制在65535我從來沒有明白。

但我仍然更喜歡向下兼容。

我也會測試@talonmies的變化。

來源

2012-05-03 00:37:44 djmj

感謝您的輸入。我會嘗試這一個。我還附加了佔用率計算器輸出。我的理解正如線程中提到的那樣正確嗎？ – Coder

即使在CUDA 4.1上的2.1設備上，最大塊尺寸仍然限制爲65535。你在哪裏看到關於CUDA 3.0 64位系列的事情？ –

sry不是64，我的意思是32位和計算能力3.0。總是把這兩個混合起來。編輯後 – djmj

如何計算塊編號

回答

相關問題