cuda-sdk的nbody代码中的线程管理

Question

当我阅读 Cuda-SDK 中的 nbody 代码时，我浏览了代码中的一些行，我发现它与他们在 GPUGems3“Fast N-Body Simulation with CUDA”中的论文有点不同。

我的问题是：首先，为什么blockIdx.x仍然涉及从全局加载内存到共享内存，如以下代码所示？

for (int tile = blockIdx.y; tile < numTiles + blockIdx.y; tile++)
{
    sharedPos[threadIdx.x+blockDim.x*threadIdx.y] =
        multithreadBodies ?
        positions[WRAP(blockIdx.x + q * tile + threadIdx.y, gridDim.x) * p + threadIdx.x] : //this line
        positions[WRAP(blockIdx.x + tile,                   gridDim.x) * p + threadIdx.x];  //this line

    __syncthreads();

    // This is the "tile_calculation" function from the GPUG3 article.
    acc = gravitation(bodyPos, acc);

    __syncthreads();
}

根据论文，它不应该是这样的吗？我想知道为什么

    sharedPos[threadIdx.x+blockDim.x*threadIdx.y] =
        multithreadBodies ?
        positions[WRAP(q * tile + threadIdx.y, gridDim.x) * p + threadIdx.x] :
        positions[WRAP(tile,                   gridDim.x) * p + threadIdx.x];

其次，在每个主体的多个线程中为什么仍然涉及threadIdx.x？它不应该是一个固定值或根本不涉及，因为总和仅归因于 threadIdx.y

if (multithreadBodies)
{
    SX_SUM(threadIdx.x, threadIdx.y).x = acc.x; //this line
    SX_SUM(threadIdx.x, threadIdx.y).y = acc.y; //this line
    SX_SUM(threadIdx.x, threadIdx.y).z = acc.z; //this line

    __syncthreads();

    // Save the result in global memory for the integration step
    if (threadIdx.y == 0)
    {
        for (int i = 1; i < blockDim.y; i++)
        {
            acc.x += SX_SUM(threadIdx.x,i).x; //this line
            acc.y += SX_SUM(threadIdx.x,i).y; //this line
            acc.z += SX_SUM(threadIdx.x,i).z; //this line
        }
    }
}

任何人都可以向我解释这一点吗？它是对更快代码的某种优化吗？

Answer 1

我是此代码和论文的作者。编号答案对应于您编号的问题。

1. 论文WRAP中没有提到blockIdx.x 对宏的偏移量，因为这是一个微优化。我什至不确定它是否值得。目的是确保不同的 SM 访问不同的 DRAM 内存库，而不是同时访问同一个内存库，以确保我们在这些负载期间最大化内存吞吐量。如果没有blockIdx.x偏移量，所有同时运行的线程块将同时访问相同的地址。由于整个算法是计算而不是带宽限制，这绝对是一个次要的优化。可悲的是，它使代码更加混乱。

2. threadIdx.y正如您所说，总和是 cross ，但每个线程都需要做一个单独的总和（每个线程计算一个单独的物体的引力）。因此我们需要使用threadIdx.x来索引（概念上的 2D）共享内存数组的右列。

在他的（不是真正正确的）答案中回答 SystmD 的问题，gridDim.y在（默认/常见）1D 块情况下只有 1。