CUDA - 沒有塊，只是未定義維度的線程

我有一些未知大小的矩陣，從兩個方向的10-20.000變化。CUDA - 沒有塊，只是未定義維度的線程

我設計了一個帶有（x; y）塊和（x; y）線程的CUDA內核。

由於矩陣的寬度/高度不是我的維數的倍數，因此讓事情起作用並且代碼變得越來越複雜以獲得聚結內存讀取變得非常可怕。

除此之外，內核越來越大，使用越來越多的寄存器來檢查正確性......所以我認爲這不是我應該採用的方式。

我的問題是：如果我完全消除塊並只創建一個x; y線程的網格會怎麼樣？ SM單位會有沒有多少塊的問題？

我可以消除塊並使用大量的線程，或者是否需要塊細分？

來源

2011-03-29 Marco A.

因爲必須將線程組織成塊，並且每個塊最多可以有512個線程，所以不能真正創建「線程網格」。但是，您可以通過每塊使用1個線程來實現此目的，這將導致1x1塊的X×Y網格。然而，這將導致非常可怕的性能，因爲以下幾個因素：

按照CUDA編程指南，一個SM可以在任何時間最多8塊的處理。這將限制你每個SM 8個線程，這不足以填充一個單一的warp。如果您擁有48個CUDA內核，則只能在任何給定時間處理384個線程。
由於SM上只有8個線程可用，因此隱藏內存延遲的變化太少。 GPU將花費大部分時間等待內存訪問完成，而不是進行任何計算。
您將無法合併內存讀取和寫入，導致較差的內存帶寬使用情況。
您將無法充分利用共享內存，因爲這是塊中線程之間的共享資源。

雖然必須確保塊中線程的正確性令人討厭，但您的性能將遠遠超過您的「線程網格」概念。

來源

2011-03-29 22:20:04 Scott

sigh ..然後我需要很多檢查代碼..謝謝 – 2011-03-29 22:22:23

@Paul：它只是一行設備代碼來檢查邊界。 – Stringer 2011-03-29 23:02:14

這不是我的情況，我正在研究nvidia SDK的合併可分卷積濾波器，順便說一句，我會看看我能做什麼 – 2011-03-30 10:47:24

下面是我用來將給定任務分成塊和網格的代碼。是的，你最終可能會啓動到很多塊（但只有很少），你可能會得到比所需的更多的實際線程，但這樣很容易和高效。查看下面的第二個代碼示例，瞭解我的簡單內核邊界檢查。 PS：我總是有block_size == 128，因爲它在多核佔用率，註冊表使用率，共享內存要求和所有內核的聚結存取之間取得了很好的平衡。

代碼來計算一個很好的網格大小（主持人）：

#define GRID_SIZE 65535 

//calculate grid size (store result in grid/block) 
void kernelUtilCalcGridSize(unsigned int num_threads, unsigned int block_size, dim3* grid, dim3* block) { 


    //block 
    block->x = block_size; 
    block->y = 1; 
    block->z = 1; 


    //number of blocks 
    unsigned int num_blocks = kernelUtilCeilDiv(num_threads, block_size); 
    unsigned int total_threads = num_blocks * block_size; 
    assert(total_threads >= num_threads); 

    //calculate grid size 
    unsigned int gy = kernelUtilCeilDiv(num_blocks, GRID_SIZE); 
    unsigned int gx = kernelUtilCeilDiv(num_blocks, gy); 
    unsigned int total_blocks = gx * gy; 
    assert(total_blocks >= num_blocks); 

    //grid 
    grid->x = gx; 
    grid->y = gy; 
    grid->z = 1; 
} 

//ceil division (rounding up) 
unsigned int kernelUtilCeilDiv(unsigned int numerator, unsigned int denominator) { 
    return (numerator + denominator - 1)/denominator; 
}

代碼來計算獨特的線程ID和檢查邊界（設備）：

//some kernel 
__global__ void kernelFoo(unsigned int num_threads, ...) { 


    //calculate unique id 
    const unsigned int thread_id = threadIdx.x; 
    const unsigned int block_id = blockIdx.x + blockIdx.y * gridDim.x; 
    const unsigned int unique_id = thread_id + block_id * blockDim.x; 


    //check range 
    if (unique_id >= num_threads) return; 

    //do the actual work 
    ... 
}

我不認爲這是一個大量努力/寄存器/代碼行來檢查正確性。

來源

2011-03-30 16:41:11 LumpN

CUDA - 沒有塊，只是未定義維度的線程

回答

相關問題