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

Posted 2023-02-16

【中文标题】cuda-sdk 的 nbody 代码中的线程管理

【英文标题】：thread management in nbody code of cuda-sdk

【发布时间】：2012-07-09 04:18:23

【问题描述】：

当我阅读 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
        
    

谁能给我解释一下？是为了更快的代码进行某种优化吗？

【参考方案1】：

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

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

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

要回答 SystmD 在他的（不是真正正确的）答案中的问题，gridDim.y 在（默认/常见）一维块案例中只有 1。

【讨论】：

马克，我认为没有 blockIdx.x 偏移 WRAP 是不需要的。您能否在for (int i = 0; i < blockDim.y; i++) 中使用0 确认SystmD 不正确

另外，为什么fractional_tile 不再处理int numTiles = n / (p * q); 中的剩余部分了

“不同的 SM 正在访问不同的 DRAM 内存块，而不是同时在同一个内存块上敲击”——后者不会利用 Fermi 等上的缓存吗？

没有分数瓦片，因为代码要求它平分。这是一个演示。我正在开发一个更好、更简单的代码版本，它只关注新 GPU 上的大问题，以消除所有这些为旧 GPU 编写的令人困惑的优化。在下一个 CUDA 版本中寻找它。

【参考方案2】：

1) 在每个块的线程同步之前（使用 __syncthreads()），数组 SharedPos 被加载到每个块（即每个图块）的共享内存中。根据算法，blockIdx.x 是图块的索引。

每个线程（索引 threadIdx.x threadIdx.y）加载共享数组 SharedPos 的一部分。 blockIdx.x 指的是 tile 的索引（没有多线程）。

2) acc是body索引blockIdx.x * blockDim.x + threadIdx.x的float3（见integrateBodies函数开头）

在与 q>4 (128 body,p =16,q=8 gridx=8) 的总和期间，我发现 multithreadBodies=true 存在一些问题。 （使用 GTX 680）。有些总和没有对整个 blockDim.y 完成......

我更改了代码以避免这种情况，它可以工作，但我不知道为什么......

if (multithreadBodies)

    SX_SUM(threadIdx.x, threadIdx.y).x = acc.x;
    SX_SUM(threadIdx.x, threadIdx.y).y = acc.y;
    SX_SUM(threadIdx.x, threadIdx.y).z = acc.z;

    __syncthreads();


        for (int i = 0; i < blockDim.y; i++) 
        
            acc.x += SX_SUM(threadIdx.x,i).x;
            acc.y += SX_SUM(threadIdx.x,i).y;
            acc.z += SX_SUM(threadIdx.x,i).z;
        

另一个问题： 在第一个循环中：

for (int tile = blockIdx.y; tile < numTiles + blockIdx.y; tile++) 

我不知道为什么从 grid.y=1 开始使用 blockIdx.y。

3) 对于更快的代码，我使用异步 H2D 和 D2D 内存副本（我的代码仅使用引力内核）。