GPGPU - CUDA：全局存儲效率

Question

我試圖找出內存寫得怎麼樣了全球訪問是合併我的內核之一的，依據「全局存儲效率」 NVidia的探查的價值（我在Fermi GPU上使用CUDA 5工具包預覽版）。

據我所知，這個值是所請求的內存事務與所執行的實際nb轉換的比率，因此反映了訪問是否都完全合併（100％效率）。現在

，爲32線程塊寬度，並考慮浮點值作爲輸入和輸出，下面的測試內核給出既爲全局負載和用於全局存儲100％的效率，如預期：

__global__ void dummyKernel(float*output,float* input,size_t pitch) 
{ 
    unsigned int x = blockIdx.x * blockDim.x + threadIdx.x; 
    unsigned int y = blockIdx.y * blockDim.y + threadIdx.y; 
    int offset = y*pitch+x; 
    float tmp = input[offset]; 
    output[offset] = tmp; 
} 

我不明白的是，爲什麼當我開始在輸入讀取和輸出寫入之間添加有用的代碼時，全局存儲效率開始下降，而我沒有改變內存寫入模式或線程塊幾何？儘管如我所料，全球加載保持100％。

難道有人請澄清爲什麼會發生這種情況嗎？我認爲，由於給定warp中的所有32個線程同時（根據定義）執行輸出存儲指令，並使用「coalescing友好」模式，我仍然可以獲得100％的任何我以前的操作，但顯然我必須誤解無論是全球商店效率的含義，還是全球商店聯合的條件。

THX，

編輯：

下面是一個例子：如果我使用這個代碼（只是增加上輸入一個 「圓」 操作），全局存儲效率從100％到95％下降

__global__ void dummyKernel(float*output,float* input,size_t pitch) 
{ 
    unsigned int x = blockIdx.x * blockDim.x + threadIdx.x; 
    unsigned int y = blockIdx.y * blockDim.y + threadIdx.y; 
    int offset = y*pitch+x; 
    float tmp = round(input[offset]); 
    output[offset] = tmp; 
} 

Answer 1

好的，對我感到羞恥，我發現了這個問題：我在調試模式下分析了這個簡單的測試代碼，它爲大多數度量標準提供了完全百搭的數字。在發佈模式下重新分析給了我預期的結果：在這兩種情況下，100％的存儲效率。

Answer 2

不確定是否是這種情況，但round可能會將其參數轉換爲double，如果存在寄存器溢出，則每個線程將訪問8個字節的內存，然後將其強制爲4個字節的tmp。訪問8個字節會將合併減少到半個扭曲。

但是，我相信註冊溢出應該不會發生，因爲內核中局部變量的數量很小。你可以用nvcc --ptxas-options = -v來檢查溢出。