0

レビュー用にドリルダウンしたコードを投稿しています。問題なくコンパイル、実行できると思いますが、関係のない部分をすべて除外したので、間違っている可能性があります。

struct Users {
    double A[96];
    double B[32];
    double C[32];
};

これは、固定長配列を使用した Users 構造です。以下に主な機能を示します。

int main(int argc, char **argv) {

    int numUsers = 10;
    Users *users = new Users[numUsers];
    double Step[96];

    for (int i = 0; i < 32; i++) {
        Step[i]      = 0.8;
        Step[i + 32] = 0.8;
        Step[i + 64] = 0.8;
    }

    for (int usr = 0; usr < numUsers; usr++) {
        for (int i = 0; i < 32; i++) {
            users[usr].A[i]      = 10;
            users[usr].A[i + 32] = 20;
            users[usr].A[i + 64] = 30;
        }
        memset(users[usr].B, 0, sizeof(double) * 32);
        memset(users[usr].C, 0, sizeof(double) * 32);
    }


    double *d_Step;
    cudaMalloc((void**)&d_Step, sizeof(double) * 96);
    cudaMemcpy(d_Step, Step, sizeof(double) * 96, cudaMemcpyHostToDevice);


    Users *deviceUsers;
    cudaMalloc((void**)&deviceUsers, sizeof(Users) * numUsers);
    cudaMemcpy(deviceUsers, users, sizeof(Users) * numUsers, cudaMemcpyHostToDevice);


    dim3 grid;
    dim3 block;

    grid.x = 1;
    grid.y = 1;
    grid.z = 1;
    block.x = 32;
    block.y = 10;
    block.z = 1;
    calc<<<grid, block >>> (deviceUsers, d_Step, numUsers);

    delete users;
    return 0;
}

ここで、ステップ配列は 96 個のビンを持つ 1D 配列であり、10 個のワープにまたがっていることに注意してください (x 方向に 32 個のスレッドがあり、私のブロックには 10 個あります)。各ワープは同じ Step 配列にアクセスします。これは、以下のカーネルで確認できます。

__global__ void calc(Users *users, double *Step, int numUsers) {

    int tId = threadIdx.x + blockIdx.x * blockDim.x;
    int uId = threadIdx.y;

    while (uId < numUsers) {

        double mean00 = users[uId].A[tId]      * Step[tId];
        double mean01 = users[uId].A[tId + 32] * Step[tId + 32];
        double mean02 = users[uId].A[tId + 64] * Step[tId + 64];

        users[uId].A[tId]      = (mean00 == 0? 0 : 1 / mean00);
        users[uId].A[tId + 32] = (mean01 == 0? 0 : 1 / mean01);
        users[uId].A[tId + 64] = (mean02 == 0? 0 : 1 / mean02);

        uId += 10;
    }
}

NVIDIA Visual Profiler を使用すると、結合された取得は 47% になりました。さらに調査したところ、各ワープがアクセスしているステップ配列がこの問題を引き起こしていることがわかりました。これを何らかの定数に置き換えると、アクセスは 100% 合体します。

Q1) 私が理解しているように、合体アクセスはバイト ラインにリンクされています。つまり、バイト ラインは、整数、ダブル バイト ラインのいずれであっても、32 の倍数でなければなりません。合体アクセスが得られないのはなぜですか?

私の知る限り、cudaはデバイスのグローバルメモリにメモリブロックを割り当てるたびに、偶数アドレスを割り当てました。したがって、始点 + 32 の場所がワープによってアクセスされる限り、アクセスは合体する必要があります。私は正しいですか？

ハードウェア

Geforce GTX 470、Compute Capability 2.0

4

1 に答える 1

1

カーネルがStepグローバル メモリから 10 回読み取りました。L1 キャッシュはグローバル メモリへの実際のアクセスを減らすことができますが、それでもプロファイラーによって非効率的なアクセス パターンとして扱われます。

私のプロファイラーはそれを「グローバル負荷効率」と名付けました。合体しているかどうかはわかりません。

于 2013-10-19T10:32:46.717 に答える