0

我有那个代码：

   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）读取Circle并检测与它的碰撞2）在改变它的速度和位置后写Circle我还有Circle的circlesConst-massive，它是由cudaMallocToSybol（）分配的。它用于检查与从 gmem 读取的主圆 C（它在寄存器中）的圆的交集。

正如我所想，我很好地使用了 const 内存，它让我获得了所有的性能：'）（我错了吗？）

当我读到对 gmem 的合并访问（是否有对其他类型的内存的合并访问？我没有找到任何有关它的信息）时，我想为我尝试一下。如您所见，Circle-structure 有 8 个类型为 float = 32 位的变量。我尝试（在代码中对其进行了注释）这样做，但是，首先，我得到了错误的答案（因为我必须不正确地从 smem 中读取，如下所述），其次，我的性能降低了 33%。为什么？我认为，它不依赖于错误的字段关系。

第二个问题，正如我在从 smem 到 C 的读取附近的代码中的注释中所写，我必须以另一种方式阅读，但如果这样做，将会有很多银行冲突，所以我会得到更少的性能...那么，我怎样才能加载没有银行冲突的合并的圈子，然后再将其写回？

ps 大小超过 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 合并阅读超过一个圈子，但我不明白，如何让它合并两个圈子..

4

1 回答 1

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]与您评论中所述的内容类似。该解决方案具有持续全局访问的优点，但使用 ank 冲突存储到 smem 中。
  2. 没有更好，因为它在读入寄存器时存在银行冲突。
  3. 由于跨步的全局内存访问，情况更糟。

回答

通过插入虚拟值可以轻松解决银行冲突。而不是使用[threadId.x * 8 + 0]你会使用[threadId.x * 9 + 0]. 请注意，您正在浪费一些共享内存（即每 9 个浮点数）来将数据分散到各个存储区。请注意，首先将数据加载到共享内存中时，您必须这样做。但是请注意，您仍然需要做很多工作才能将这些Circle结构加载到那里。这导致我

更好的答案

只是不要Circle在全局内存中使用结构数组。通过使用多个浮点数组来反转您的内存模式。一个用于 a 的每个组件Circle。然后，您可以直接加载到寄存器中。

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

根本没有更多的共享内存，由于更少的指针计算，更少的寄存器，连续的全局访问模式，没有银行冲突。全部免费！

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

于 2013-12-21T10:31:54.453 回答