解決衝突 - 嘗試合併GMEM訪問，使用SMEM，但銀行衝突

Question

我有代碼：

struct __declspec(align(32)) Circle 
{ 
    float x, y; 
    float prevX, prevY; 
    float speedX, speedY; 
    float mass; 
    float radius; 

void init(const int _x, const int _y, const float _speedX = 0.0f, const float _speedY = 0.0f, 
    const float _radius = CIRCLE_RADIUS_DEFAULT, 
    const float _mass = CIRCLE_MASS_DEFAULT); 
}; 

，第二個：

/*smem[threadIdx.x] = *(((float*)cOut) + threadIdx.x); 
     smem[threadIdx.x + blockDim.x] = *(((float*)cOut) + threadIdx.x + blockDim.x); 
     smem[threadIdx.x + blockDim.x * 2] = *(((float*)cOut) + threadIdx.x + blockDim.x * 2); 
     smem[threadIdx.x + blockDim.x * 3] = *(((float*)cOut) + threadIdx.x + blockDim.x * 3); 
     smem[threadIdx.x + blockDim.x * 4] = *(((float*)cOut) + threadIdx.x + blockDim.x * 4); 
     smem[threadIdx.x + blockDim.x * 5] = *(((float*)cOut) + threadIdx.x + blockDim.x * 5); 
     smem[threadIdx.x + blockDim.x * 6] = *(((float*)cOut) + threadIdx.x + blockDim.x * 6); 
     smem[threadIdx.x + blockDim.x * 7] = *(((float*)cOut) + threadIdx.x + blockDim.x * 7);*/ 
     __syncthreads(); 
     /*float x, y; 
     float prevX, prevY; 
     float speedX, speedY; 
     float mass; 
     float radius;*/ 
     /*c.x = smem[threadIdx.x]; 
     c.y = smem[threadIdx.x + blockDim.x]; //there must be [threadId.x * 8 + 0] 
     c.prevX = smem[threadIdx.x + blockDim.x * 2]; //[threadId.x * 8 + 1] and e.t.c. 
     c.prevY = smem[threadIdx.x + blockDim.x * 3]; 
     c.speedX = smem[threadIdx.x + blockDim.x * 4]; 
     c.speedY = smem[threadIdx.x + blockDim.x * 5]; 
     c.mass = smem[threadIdx.x + blockDim.x * 6]; 
     c.radius = smem[threadIdx.x + blockDim.x * 7];*/ 
     c = cOut[j]; 
     //c = *((Circle*)(smem + threadIdx * SMEM)); 

有2 GMEM（我的意思是全球內存）訪問： 1）讀取圓並檢測它與它的碰撞 2）在改變它的速度和位置後寫入圓 另外我還有circleConst-mass的Circle，它是由cudaMallocToSybol（）分配的。它用於檢查從gmem讀取的主圓C（它在寄存器中）的圓的交集。

當我想到，我使用const，記憶好，它獲得了我所有的性能：'）（我錯了）

當我讀到凝聚的訪問GMEM（有合併獲得其他類型的？記憶？我沒有找到任何有關它的信息），我想爲我嘗試它。如你所見，Circle-structure有8個變量，類型爲float = 32位。我嘗試過（在代碼中評論）做這件事，但是，首先，我得到了一個錯誤的答案（因爲我必須從不正確的地方讀到，下面會提到），其次，我的表現會降低33％。爲什麼？我認爲，這不取決於錯誤的領域關係。

而第二個問題，正如我在代碼的評論中寫到從smem到C的附近的代碼所寫的，我必須讀另一種方式，但如果我這樣做，會有很多銀行衝突，所以我會獲得更少的性能...... 那麼，我怎樣才能加載圈沒有銀行衝突coalasced，然後，把它寫回來？

p.s大小超過4 * float的結構是否位於寄存器中？

---

更新： 最新的版本是：

#define CF (9) //9 because the primary struct has 8 floats, so 1 is for wasting 

i = blockIdx.x * blockDim.x; 
     smem[threadIdx.x + blockDim.x * 0 + blockDim.x * 0/(CF - 1) + threadIdx.x/(CF - 1)] = *(((float*)(cOut + i)) + threadIdx.x + blockDim.x * 0); 
     smem[threadIdx.x + blockDim.x * 1 + blockDim.x * 1/(CF - 1) + threadIdx.x/(CF - 1)] = *(((float*)(cOut + i)) + threadIdx.x + blockDim.x * 1); 
     smem[threadIdx.x + blockDim.x * 2 + blockDim.x * 2/(CF - 1) + threadIdx.x/(CF - 1)] = *(((float*)(cOut + i)) + threadIdx.x + blockDim.x * 2); 
     smem[threadIdx.x + blockDim.x * 3 + blockDim.x * 3/(CF - 1) + threadIdx.x/(CF - 1)] = *(((float*)(cOut + i)) + threadIdx.x + blockDim.x * 3); 
     smem[threadIdx.x + blockDim.x * 4 + blockDim.x * 4/(CF - 1) + threadIdx.x/(CF - 1)] = *(((float*)(cOut + i)) + threadIdx.x + blockDim.x * 4); 
     smem[threadIdx.x + blockDim.x * 5 + blockDim.x * 5/(CF - 1) + threadIdx.x/(CF - 1)] = *(((float*)(cOut + i)) + threadIdx.x + blockDim.x * 5); 
     smem[threadIdx.x + blockDim.x * 6 + blockDim.x * 6/(CF - 1) + threadIdx.x/(CF - 1)] = *(((float*)(cOut + i)) + threadIdx.x + blockDim.x * 6); 
     smem[threadIdx.x + blockDim.x * 7 + blockDim.x * 7/(CF - 1) + threadIdx.x/(CF - 1)] = *(((float*)(cOut + i)) + threadIdx.x + blockDim.x * 7); 

c.x =  smem[threadIdx.x * CF + 0]; 
    c.y =  smem[threadIdx.x * CF + 1]; 
    c.prevX = smem[threadIdx.x * CF + 2]; 
    c.prevY = smem[threadIdx.x * CF + 3]; 
    c.speedX = smem[threadIdx.x * CF + 4]; 
    c.speedY = smem[threadIdx.x * CF + 5]; 
    c.mass = smem[threadIdx.x * CF + 6]; 
    c.radius = smem[threadIdx.x * CF + 7]; 

是不是正確的方式使用存取權限合併SMEM GMEM？我的意思是，我害怕BlockDim.x * 1/(CF - 1) + threadIdx.x/(CF - 1)。 我想，我沒有得到一些提升，因爲它不允許gmem合併讀取超過一個圓，但我不明白，如何使它合併兩個圓..

Answer 1

免責聲明

請注意，此答案包含更多的問題比答案。另外請注意，我猜測很多，因爲我沒有得到你的問題和源代碼的大部分。

重建

所以我猜你的全局內存是Circle結構數組。您似乎已經通過將它們的每個浮點數分別加載到共享內存中來優化加載這些圓圈。這樣你可以獲得連續的訪問模式，而不是跨步式訪問模式。我還在這裏嗎？

所以，現在你已經裝載blockDim.x圈到共享內存中合作，你想從它讀了一圈c爲每個線程，你似乎已經嘗試了3種不同的方式：

1. 裝載c從跨進共享內存
   （c.prevX = smem[threadIdx.x + blockDim.x * 2];等）
2. 裝載c直接從共享存儲器
   （c = *((Circle*)(smem + threadIdx * SMEM));）
3. 裝載c直接從全局內存
   （c = cOut[j];）

仍然是正確的？

評價

1. 當你加載圈子到像我以前所描述的方式共享內存沒有任何意義。所以你可能試過了不同的加載模式。在您的評論中提到的[threadId.x * 8 + 0]的內容。該解決方案具有持續全局訪問的優點，但使用腳本衝突存儲在屏幕上。
2. 沒有更好的，因爲它讀成寄存器時有庫衝突。
3. 是因爲跨入全球內存訪問更糟。

回答

銀行的衝突很容易被插入虛擬值解決。而不是使用[threadId.x * 8 + 0]的你會使用[threadId.x * 9 + 0]。請注意，您正在浪費一些共享內存（即每第九個浮點）來分散跨銀行的數據。請注意，首先將數據加載到共享內存時，您必須執行相同的操作。但請注意，您仍然在做很多工作來將這些Circle結構加載到那裏。這使我的

甚至更​​好的答案

只要不使用的全局內存Circle結構數組。改爲使用多個float數組來反轉你的記憶模式。一個用於Circle的每個組件。然後您可以直接加載到寄存器中。

c.x = gmem_x[j]; 
c.y = gmem_y[j]; 
... 

沒有更多的共享內存在所有，更少的寄存器由於指針計算更少，連續全球訪問模式，沒有組衝突。所有這一切都是免費的！

現在你可能認爲準備數據時，主機上獲得結果後面有一個缺點吧。我最好的（也是最後的）猜測是，它總體上還是要快得多，因爲你可能會在每一幀都啓動內核，並使用着色器進行可視化，而無需將數據傳回主機或連續多次啓動內核在下載結果之前。正確？