2

我正在使用 CUDA 对可能很大的 3D 数据集进行计算。我认为最好先看一个简短的代码片段：

void launch_kernel(/*arguments . . . */){
    int bx = xend-xstart, by = yend-ystart, bz = zend-zstart;

    dim3 blocks(/*dimensions*/);
    dim3 threads(/*dimensions*/);
    kernel<<blocks, threads>>();
}

我有一组 3D 单元，我需要启动一个内核来计算每个单元。问题是输入大小可能超出 GPU 的能力，特别是线程。所以这样的代码：

void launch_kernel(/*arguments . . . */){
       int bx = xend-xstart, by = yend-ystart, bz = zend-zstart;

       dim3 blocks(bx,by,1);
       dim3 threads(bz);
       kernel<<blocks, threads>>();
   }

……不好用。因为如果尺寸是 1000x1000x1000 怎么办？- 我无法在每个块中启动 1000 个线程。或者更好的是，如果尺寸是 5x5x1000 怎么办？- 现在我几乎没有启动任何块，但内核需要启动 5x5x512 b/c 的硬件，每个线程将执行 2 次计算。我也不能只是混搭我的所有维度，将一些 z 放在块中，一些放在线程 b/c 中，我需要能够识别内核中的维度。目前：

__global__ void kernel(/*arguments*/){
    int x = xstart + blockIdx.x;
    int y = ystart + blockIdx.y;
    int z = zstart + threadIdx.x;
    if(x < xend && y < yend && z < zend){
        //calculate
    }
}

我需要一种可靠、有效的方法来计算这些变量：

块 x 维度，块 y 维度，线程 x（以及 y? 和 z?），一旦我通过 blockIdx 和 threadIdx 在内核中的 x,y,z，如果输入超过硬件，a 的数量我在内核计算中的 for 循环中为每个维度采取“步骤”。

如果您有任何问题，请询问。这是一个难题，一直困扰着我（特别是因为我启动的块/线程的数量是性能的主要组成部分）。该代码需要在其针对不同数据集的决策中自动化，我不确定如何有效地做到这一点。先感谢您。

4

2 回答 2

3

我认为你在这里把事情复杂化了。基本问题似乎是您需要在 1000 x 1000 x 1000 计算域上运行内核。因此，您需要 1000000000 个线程，这完全在所有 CUDA 兼容硬件的能力范围内。因此，只需使用标准的 2D CUDA 执行网格，其中至少包含进行计算所需的线程数（如果您不明白如何做到这一点，请发表评论，我会将其添加到答案中），然后在您的内核调用中一个小的设置功能是这样的：

__device__ dim3 thread3d(const int dimx, const int dimxy)
{
    // The dimensions of the logical computational domain are (dimx,dimy,dimz)
    // and dimxy = dimx * dimy
    int tidx = threadIdx.x + blockIdx.x * blockDim.x;
    int tidy = threadIdx.y + blockIdx.y * blockDim.y;
    int tidxy = tidx + gridDim.x * tidy;

    dim3 id3d;
    id3d.z = tidxy / dimxy;
    id3d.y = tidxy / (id3d.z * dimxy);
    id3d.x = tidxy - (id3d.z * dimxy - id3d.y * dimx);

    return id3d;
}

[免责声明：用浏览器编写，从未编译，从未运行，从未测试。使用风险自负]。

此函数将从 CUDA 2D 执行网格返回 3D 域 (dimx,dimy,dimz) 中的“逻辑”线程坐标。在内核的开头调用它，如下所示：

__global__ void kernel(arglist, const int dimx, const int dimxy)
{
    dim3 tid = thread3d(dimx, dimxy);

    // tid.{xyx} now contain unique 3D coordinates on the (dimx,dimy,dimz) domain
    .....
}

请注意，在设置该网格时存在大量整数计算开销，因此您可能需要考虑为什么您真的需要 3D 网格。您会惊讶于它实际上没有必要的次数，并且可以避免大部分设置开销。

于 2012-06-08T06:24:38.487 回答
1

我将首先用于cudaGetDeviceProperties()查找您的 GPU 的计算能力，以便您确切地知道您的 GPU 允许每个块有多少线程（如果您的程序需要通用化，以便它可以在任何支持 CUDA 的设备上运行）。

然后，使用这个数字，我会做一个大的嵌套if语句来测试你输入的维度。如果所有维度都足够小，则可以拥有一组 (bx,by,bz) 线程（不太可能）。如果这不起作用，则找到可以放入一个块中的最大维度（或两个维度）并据此进行分区。如果这不起作用，那么您将不得不对最小维度进行分区，以便将其中的一些块放入一个块中 - 例如假设和的(MAX_NUMBER_THREADS_PER_BLOCK,1,1)线程和块。(bx/MAX_NUMBER)THREADS_PER_BLOCK,by,bz)bx<by<bzbx>MAX_NUMBER_THREADS_PER_BLOCK

对于每种情况，您都需要不同的内核，这有点痛苦，但归根结底，这是一项可行的工作。

于 2012-06-08T00:37:19.623 回答