3

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

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

4

2 回答 2

5

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

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

  2. The sum is across threadIdx.y, as you say, but each thread needs to do a separate sum (each thread computes gravitation for a separate body). Therefore we need to use threadIdx.x to index the right column of the (conceptually 2D) shared memory array.

To Answer SystmD's question in his (not really correct) answer, gridDim.y is only 1 in the (default/common) 1D block case.

于 2012-09-12T02:04:58.393 回答
0

1)数组SharedPos在每个block的线程同步之前加载到每个block(即每个tile)的共享内存中(用__syncthreads())。根据算法,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 内存副本(我的代码仅使用引力内核)。

于 2012-07-10T09:19:57.110 回答