0

我正在发布一个深入研究的代码以供审查。我相信它应该可以毫无问题地编译和执行，但由于我排除了所有不相关的部分，我可能犯了一些错误。

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

这是我的带有固定长度数组的用户结构。下面给出主要功能。

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

请注意，Step 数组是具有 96 个 bin 的 1D 数组，我跨越 10 个 warp（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%。我进一步调查并发现每个经线正在访问的 Step 数组会导致此问题。如果我用一些常量替换它，则访问是 100% 合并的。

Q1）据我了解，合并访问链接到字节行，即字节行必须是 32 的倍数，无论它们是整数还是双字节行。为什么我没有获得合并访问？

据我所知，cuda 每当在设备全局内存中分配一个内存块时，它都会为其分配一个偶数地址。因此，只要起点 + 32 位置由经线访问，就应该合并访问。我对么？

硬件

Geforce GTX 470，计算能力 2.0

4

1 回答 1

1

你的内核Step从全局内存中读取了 10 次。虽然 L1 缓存可以减少对 global mem 的实际访问，但它仍然被 profiler 视为低效的访问模式。

我的分析器将其命名为“全局负载效率”。它没有说它是否合并。

于 2013-10-19T10:32:46.717 回答