不連貫的賣場

爲什麼這個內核產生不連貫的賣場不連貫的賣場

__global__ void reverseArrayBlock(int *d_out, int *d_in) 
{ 
    int inOffset = blockDim.x * blockIdx.x; 
    int outOffset = blockDim.x * (gridDim.x - 1 - blockIdx.x); 
    int in = inOffset + threadIdx.x; 
    int out = outOffset + (blockDim.x - 1 - threadIdx.x); 
    d_out[out] = d_in[in]; 
}

，這一次不

__global__ void reverseArrayBlock(int *d_out, int *d_in) 
{ 
    extern __shared__ int s_data[]; 

    int inOffset = blockDim.x * blockIdx.x; 
    int in = inOffset + threadIdx.x; 

    // Load one element per thread from device memory and store it 
    // *in reversed order* into temporary shared memory 
    s_data[blockDim.x - 1 - threadIdx.x] = d_in[in]; 

    // Block until all threads in the block have written their data to shared mem 
    __syncthreads(); 

    // write the data from shared memory in forward order, 
    // but to the reversed block offset as before 

    int outOffset = blockDim.x * (gridDim.x - 1 - blockIdx.x); 

    int out = outOffset + threadIdx.x; 
    d_out[out] = s_data[threadIdx.x]; 
}

我知道，第二個是使用共享內存。但是當我看到d_out的指示時，它們在內核中似乎是相同的。你能幫我理解嗎？

來源

2011-06-02 WannabeCoder

合併要求地址在變形內遵循「基本+ tid」模式，其中tid是線索索引的縮寫。換句話說，隨着tid的增加，地址也會增加。你的評論稱之爲「遠期訂單」。在第一個內核中，生成地址，使得隨着tid增加，地址減少，即訪問處於「反向順序」。

來源

2011-06-02 18:48:08 njuffa

感謝您的澄清。 – WannabeCoder 2011-06-02 21:26:46

請注意，示例代碼只能在Compute Capability 1.1和早期設備上生成不連貫的商店。從1.2設備開始，此訪問模式完全合併/一致。 – harrism 2011-06-06 09:19:53

在我們開始之前，您需要了解寫入共享內存比寫入全局內存要便宜得多。

考慮到這一點，讓我們說，我們是反相的陣列1-> 32

方法一個執行此：在寫入線程1從位置x讀取，線程2從讀取（X + 1），線程3從位置（x + 2）讀取...線程32從位置（x + 31）讀取。

您可以在2（如果對齊）或3（如果未對齊）讀取中讀取整個內存塊，因爲操作是以半塊（16個線程）塊完成的。寫入線程1寫入位置（y + 31），線程2寫入（y + 30），線程3寫入位置（y + 29）...線程32寫入位置（y）。

儘管它們正在寫入連續的內存塊，但它們的順序相反。除非你使用了一些最新的硬件（即使有它，我也很懷疑），但是需要32次寫入才能執行操作。

至於第二種情況，你正在做32個反向寫入共享內存和32個反向讀取共享內存，這並不昂貴。

現在，您已經按相反順序讀取數據，您可以按正確的順序寫入全局內存。

線程1寫入位置y，線程2寫入位置y + 1等等。

底線，您節省了執行32（方法1）-3（方法2）= 29寫入全局內存所需的時間。

來源

2011-06-02 20:43:15

不連貫的賣場

回答

相關問題