解决冲突 - 尝试合并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]; 
... 

没有更多的共享内存在所有，更少的寄存器由于指针计算更少，连续全球访问模式，没有组冲突。所有这一切都是免费的！

现在你可能认为准备数据时，主机上获得结果后面有一个缺点吧。我最好的（也是最后的）猜测是，它总体上还是要快得多，因为你可能会在每一帧都启动内核，并使用着色器进行可视化，而无需将数据传回主机或连续多次启动内核在下载结果之前。正确？