cuda - 衝突の解決 - smem を使用して gmem アクセスを合体させようとしますが、銀行の競合が発生します

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);
};

そして2番目のもの：

/*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を書き込みます gmem から読み込まれた主円 C (レジスタにある) の円との交差をチェックするために使用されます。

私が思うに、私は const メモリをうまく使って、そのパフォーマンスをすべて得ています :') (私は間違っていますか?)

gmem への合体アクセスについて読んだとき (他のタイプのメモリへの合体アクセスはありますか? それに関する情報は見つかりませんでした)、試してみたかったのです。ご覧のとおり、Circle-structure には float = 32bits 型の 8 つの変数があります。私はそれをやろうとしましたが（コードではコメントされています）、最初に間違った答えが得られました（後述のsme​​mから正しく読み取らなければならないため）、次にパフォーマンスが33％低下します。なんで？間違ったフィールドの関係には依存しないと思います。

そして2つ目の質問ですが、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 が複数のサークルを結合して読み取ることができないため、ブーストが得られなかったと思いますが、2 つのサークルを結合する方法がわかりません..