如何避免將數據從全局數據加載到共享內存時發生銀行衝突

存在於計算能力1.3 GPU的全局內存中的一個無符號字符數組的步進存取問題。爲了繞過全局存儲器的聚結的要求，螺紋依次訪問全局存儲器和複製僅使用2存儲器事務以下示例陣列到共享存儲器：如何避免將數據從全局數據加載到共享內存時發生銀行衝突

#include <cuda.h> 
#include <stdio.h> 
#include <stdlib.h> 

__global__ void kernel (unsigned char *d_text, unsigned char *d_out) { 

    int idx = blockIdx.x * blockDim.x + threadIdx.x; 

    extern __shared__ unsigned char s_array[]; 

    uint4 *uint4_text = (uint4 *) d_text; 
    uint4 var; 

    //memory transaction 
    var = uint4_text[0]; 

    uchar4 c0 = *reinterpret_cast<uchar4 *>(&var.x); 
    uchar4 c4 = *reinterpret_cast<uchar4 *>(&var.y); 
    uchar4 c8 = *reinterpret_cast<uchar4 *>(&var.z); 
    uchar4 c12 = *reinterpret_cast<uchar4 *>(&var.w); 

    s_array[threadIdx.x*16 + 0] = c0.x; 
    s_array[threadIdx.x*16 + 1] = c0.y; 
    s_array[threadIdx.x*16 + 2] = c0.z; 
    s_array[threadIdx.x*16 + 3] = c0.w; 

    s_array[threadIdx.x*16 + 4] = c4.x; 
    s_array[threadIdx.x*16 + 5] = c4.y; 
    s_array[threadIdx.x*16 + 6] = c4.z; 
    s_array[threadIdx.x*16 + 7] = c4.w; 

    s_array[threadIdx.x*16 + 8] = c8.x; 
    s_array[threadIdx.x*16 + 9] = c8.y; 
    s_array[threadIdx.x*16 + 10] = c8.z; 
    s_array[threadIdx.x*16 + 11] = c8.w; 

    s_array[threadIdx.x*16 + 12] = c12.x; 
    s_array[threadIdx.x*16 + 13] = c12.y; 
    s_array[threadIdx.x*16 + 14] = c12.z; 
    s_array[threadIdx.x*16 + 15] = c12.w; 

    d_out[idx] = s_array[threadIdx.x*16]; 
} 

int main (void) { 

    unsigned char *d_text, *d_out; 

    unsigned char *h_out = (unsigned char *) malloc (32 * sizeof (unsigned char)); 
    unsigned char *h_text = (unsigned char *) malloc (32 * sizeof (unsigned char)); 

    int i; 

    for (i = 0; i < 32; i++) 
     h_text[i] = 65 + i; 

    cudaMalloc ((void**) &d_text, 32 * sizeof (unsigned char)); 
    cudaMalloc ((void**) &d_out, 32 * sizeof (unsigned char)); 

    cudaMemcpy (d_text, h_text, 32 * sizeof (unsigned char), cudaMemcpyHostToDevice); 

    kernel<<<1,32,16128>>>(d_text, d_out); 

    cudaMemcpy (h_out, d_out, 32 * sizeof (unsigned char), cudaMemcpyDeviceToHost); 

    for (i = 0; i < 32; i++) 
     printf("%c\n", h_out[i]); 

    return 0; 
}

的問題是，組衝突在將數據複製到共享內存時發生（由nvprof報告，上述示例的衝突爲384次），這會導致線程的序列化訪問。

共享內存分爲16個（或更新設備體系結構中的32個）32位存儲區，以便同時服務相同半經線的16個線程。數據交錯存儲在第i個32位字始終存儲在i％16-1共享存儲區中。

由於每個線程讀取一個內存事務的16個字節，這些字符將以交叉方式存儲到共享內存中。這會導致線程0,4,8,12之間的衝突; 1，5，9，13; 2，6，10，14; 3，7，11，15是同一個半經線。消除體衝突一個天真的方法是使用的if/else分支將數據存儲在類似以下內容的循環方式共享內存，但導致一些嚴重的線程分歧：

int tid16 = threadIdx.x % 16; 

if (tid16 < 4) { 

    s_array[threadIdx.x * 16 + 0] = c0.x; 
    s_array[threadIdx.x * 16 + 1] = c0.y; 
    s_array[threadIdx.x * 16 + 2] = c0.z; 
    s_array[threadIdx.x * 16 + 3] = c0.w; 

    s_array[threadIdx.x * 16 + 4] = c4.x; 
    s_array[threadIdx.x * 16 + 5] = c4.y; 
    s_array[threadIdx.x * 16 + 6] = c4.z; 
    s_array[threadIdx.x * 16 + 7] = c4.w; 

    s_array[threadIdx.x * 16 + 8] = c8.x; 
    s_array[threadIdx.x * 16 + 9] = c8.y; 
    s_array[threadIdx.x * 16 + 10] = c8.z; 
    s_array[threadIdx.x * 16 + 11] = c8.w; 

    s_array[threadIdx.x * 16 + 12] = c12.x; 
    s_array[threadIdx.x * 16 + 13] = c12.y; 
    s_array[threadIdx.x * 16 + 14] = c12.z; 
    s_array[threadIdx.x * 16 + 15] = c12.w; 

} else if (tid16 < 8) { 

    s_array[threadIdx.x * 16 + 4] = c4.x; 
    s_array[threadIdx.x * 16 + 5] = c4.y; 
    s_array[threadIdx.x * 16 + 6] = c4.z; 
    s_array[threadIdx.x * 16 + 7] = c4.w; 

    s_array[threadIdx.x * 16 + 8] = c8.x; 
    s_array[threadIdx.x * 16 + 9] = c8.y; 
    s_array[threadIdx.x * 16 + 10] = c8.z; 
    s_array[threadIdx.x * 16 + 11] = c8.w; 

    s_array[threadIdx.x * 16 + 12] = c12.x; 
    s_array[threadIdx.x * 16 + 13] = c12.y; 
    s_array[threadIdx.x * 16 + 14] = c12.z; 
    s_array[threadIdx.x * 16 + 15] = c12.w; 

    s_array[threadIdx.x * 16 + 0] = c0.x; 
    s_array[threadIdx.x * 16 + 1] = c0.y; 
    s_array[threadIdx.x * 16 + 2] = c0.z; 
    s_array[threadIdx.x * 16 + 3] = c0.w; 

} else if (tid16 < 12) { 

    s_array[threadIdx.x * 16 + 8] = c8.x; 
    s_array[threadIdx.x * 16 + 9] = c8.y; 
    s_array[threadIdx.x * 16 + 10] = c8.z; 
    s_array[threadIdx.x * 16 + 11] = c8.w; 

    s_array[threadIdx.x * 16 + 12] = c12.x; 
    s_array[threadIdx.x * 16 + 13] = c12.y; 
    s_array[threadIdx.x * 16 + 14] = c12.z; 
    s_array[threadIdx.x * 16 + 15] = c12.w; 

    s_array[threadIdx.x * 16 + 0] = c0.x; 
    s_array[threadIdx.x * 16 + 1] = c0.y; 
    s_array[threadIdx.x * 16 + 2] = c0.z; 
    s_array[threadIdx.x * 16 + 3] = c0.w; 

    s_array[threadIdx.x * 16 + 4] = c4.x; 
    s_array[threadIdx.x * 16 + 5] = c4.y; 
    s_array[threadIdx.x * 16 + 6] = c4.z; 
    s_array[threadIdx.x * 16 + 7] = c4.w; 

} else { 

    s_array[threadIdx.x * 16 + 12] = c12.x; 
    s_array[threadIdx.x * 16 + 13] = c12.y; 
    s_array[threadIdx.x * 16 + 14] = c12.z; 
    s_array[threadIdx.x * 16 + 15] = c12.w; 

    s_array[threadIdx.x * 16 + 0] = c0.x; 
    s_array[threadIdx.x * 16 + 1] = c0.y; 
    s_array[threadIdx.x * 16 + 2] = c0.z; 
    s_array[threadIdx.x * 16 + 3] = c0.w; 

    s_array[threadIdx.x * 16 + 4] = c4.x; 
    s_array[threadIdx.x * 16 + 5] = c4.y; 
    s_array[threadIdx.x * 16 + 6] = c4.z; 
    s_array[threadIdx.x * 16 + 7] = c4.w; 

    s_array[threadIdx.x * 16 + 8] = c8.x; 
    s_array[threadIdx.x * 16 + 9] = c8.y; 
    s_array[threadIdx.x * 16 + 10] = c8.z; 
    s_array[threadIdx.x * 16 + 11] = c8.w; 
}

任何人都可以想出一個更好的解決方案呢？我已經研究過SDK的縮減例子，但我不確定它適用於我的問題。

來源

2012-11-01 charis

我認爲DWORD複製比單字節複製要快。試試這個你的榜樣，而不是：

for(int i = 0; i < 4; i++) 
{ 
    ((int*)s_array)[4 * threadIdx.x + i] = ((int*)d_text)[i]; 
}

來源

2012-11-01 23:51:56 otter

我實際上將uint4字（128位）直接存儲到共享內存中。每字節應對是嘗試找到解決銀行衝突問題的臨時方法 – charis

授予的代碼將導致銀行的衝突，但是，這並不意味着它是任何慢。

在您的計算能力1.3 GPU上，具有雙向銀行衝突的共享內存事務只需要比沒有銀行衝突的事務多兩個週期。在兩個週期內，您甚至無法執行單個指令來解決銀行衝突。與無衝突訪問相比，4路銀行衝突使用6個週期，這足以執行一次額外的無衝突共享內存訪問。在你的情況下，代碼很可能受全局內存帶寬（和延遲，這是幾百個週期，即比我們在這裏討論的2..6個週期大兩個數量級）的限制。所以，如果SM只是空閒等待全局內存中的數據，那麼您可能會有足夠的備用週期。然後銀行衝突可以使用這些週期，而不會減慢您的代碼的所有。

確保編譯器將.x，.y，.z和.w的四個按字節存儲合併到一個32位訪問中將更爲重要。使用cuobjdump -sass查看編譯後的代碼，看看是否是這種情況。如果不是，請按照Otter的建議來改用字轉移。

如果您只是從d_text中讀取數據，而不是從內核中讀取數據，那麼您也可以使用它的紋理，它仍然會比內存衝突的內核慢，但可能會提供其他優勢來提高整體速度（例如，如果無法保證全局內存中數據的正確對齊）。

另一方面，您的替代銀行衝突免費代碼將快速的256字節全局內存拆分爲四個64位事務，這些事務的效率會低很多，並且可能會超出正在運行的最大內存事務數量你會產生全面的四百到幾千個全局內存延遲週期。
爲避免這種情況，您需要首先使用256字節寬的讀取操作將數據傳輸到寄存器，然後以無衝突銀行衝突的方式將數據從寄存器移入共享內存。不過，只有register-> shmem移動的代碼將比我們試圖解決的六個週期佔用更多的代碼。

來源

2012-11-02 02:24:07 tera

從內存中提取單詞後發生銀行衝突。我不確定全局內存延遲是否可以有效地隱藏它們。每字節複製僅用於測試目的。我實際上是直接將uint4字複製到共享內存中。 d_text的每個字節只讀取一次，所以AFAIK紋理緩存不會有任何好處。對於大小爲116.234.496字節的d_text數組，將它們存儲在共享內存中，然後從共享內存中讀取以處理它們將導致3736445個衝突。 SDK中的縮減示例要求在避免銀行衝突時提高2倍以上 – charis

因此，銀行衝突將在GTX 260上消耗大約0.25毫秒的實時時間，而116.234.496字節的讀取時間大約需要1毫秒。你的內核需要多長時間才能執行？不過我的主要觀點是不同的：硬件已經採取瞭解決bank衝突的最佳方法，您不能編寫代碼來執行相同的任務，即每個線程將16個連續字節傳輸到共享內存的速度更快。你唯一能做的就是重新排列數據或算法的佈局，但爲了解決這個問題，我們需要更多關於你想要實現的信息。 – tera

例如，如果您可以在寫入共享內存之前將每個線程的16個字節處理爲寄存器內的四個字節，那麼這將完全消除存儲庫衝突（但如果必須進行此處理，則只會提供加速）。在可能的情況下，減少案例是一個理想的例子。 – tera

如何避免將數據從全局數據加載到共享內存時發生銀行衝突

回答

相關問題