11

最近我一直在 CUDA 上进行字符串比较工作，我想知道 __global__ 函数如何在找到我正在寻找的确切字符串时返回一个值。

我的意思是，我需要包含大量线程的 __global__ 函数来同时在一个大字符串池中找到某个字符串，我希望一旦捕获到确切的字符串，__global__ 函数可以停止所有线程并返回回到主要功能，并告诉我“他做到了”！

我正在使用 CUDA C。我怎么可能做到这一点？

4

3 回答 3

21

在 CUDA（或 NVIDIA GPU）中，一个线程无法中断所有正在运行的线程的执行。一旦找到结果，您就不能立即退出内核，这在今天是不可能的。

但是你可以让所有线程在一个线程找到结果后尽快退出。这是您将如何做到这一点的模型。

__global___ void kernel(volatile bool *found, ...) 
{
    while (!(*found) && workLeftToDo()) {

       bool iFoundIt = do_some_work(...); // see notes below

       if (iFoundIt) *found = true;
    }
}

对此的一些说明。

  1. 注意使用volatile. 这个很重要。
  2. 确保在启动内核之前初始化found（必须是设备指针） ！false
  3. 当另一个线程更新时，线程不会立即退出found。它们只会在下次返回到 while 循环的顶部时退出。
  4. 你如何实施do_some_work很重要。如果工作量太大（或变量太大），则找到结果后退出的延迟会很长（或变量）。如果工作太少，那么您的线程将花费大部分时间检查found而不是做有用的工作。
  5. do_some_work还负责分配任务（即计算/递增索引），以及如何做到这一点是特定于问题的。
  6. 如果您启动的块数远大于当前 GPU 上内核的最大占用率，并且在第一次运行的线程块“波”中未找到匹配项，则此内核（以及下面的内核）可能会死锁. 如果在第一波中找到匹配项，则后面的块将仅在 之后运行found == true，这意味着它们将启动，然后立即退出。解决方案是同时启动尽可能多的块（也称为“最大启动”），并相应地更新您的任务分配。
  7. 如果任务数量比较少，可以while用 an代替if，运行刚好够覆盖任务数量的线程。那么就没有死锁的机会（但前一点的第一部分适用）。
  8. workLeftToDo()是特定于问题的，但是当没有工作要做时它会返回 false ，这样我们就不会在找不到匹配项的情况下陷入僵局。

现在，上述情况可能会导致过度的分区驻留（所有线程都在同一内存上），尤其是在没有 L1 缓存的旧架构上。因此，您可能想要编写一个稍微复杂一点的版本，使用每个块的共享状态。

__global___ void kernel(volatile bool *found, ...) 
{
    volatile __shared__ bool someoneFoundIt;

    // initialize shared status
    if (threadIdx.x == 0) someoneFoundIt = *found;
    __syncthreads();

    while(!someoneFoundIt && workLeftToDo()) {

       bool iFoundIt = do_some_work(...); 

       // if I found it, tell everyone they can exit
       if (iFoundIt) { someoneFoundIt = true; *found = true; }

       // if someone in another block found it, tell 
       // everyone in my block they can exit
       if (threadIdx.x == 0 && *found) someoneFoundIt = true;

       __syncthreads();
    }
}

这样，每个块有一个线程轮询全局变量，并且只有找到匹配项的线程才会写入它，因此全局内存流量被最小化。

另外：__global__ 函数是无效的，因为很难定义如何将值从 1000 个线程返回到单个 CPU 线程。用户在设备或零拷贝内存中设计适合其目的的返回数组是微不足道的，但很难制作通用机制。

免责声明：在浏览器中编写的代码，未经测试，未经验证。

于 2012-09-20T04:22:56.240 回答
5

如果您喜欢冒险，停止内核执行的另一种方法是只执行

// (write result to memory here)
__threadfence();
asm("trap;");

如果找到答案。

这不需要轮询内存，但不如 Mark Harris 建议的解决方案，因为它使内核退出并出现错误情况。这可能会掩盖实际错误（因此请务必以清楚地允许从错误中区分成功执行的方式写出您的结果），并且可能会导致其他问题或降低整体性能，因为驱动程序将此视为异常。

如果您正在寻找一个安全且简单的解决方案，请改用 Mark Harris 的建议。

于 2012-09-20T11:09:10.883 回答
0

全局函数并不像您想象的那样真正包含大量线程。它只是一个内核，在设备上运行的函数，通过传递指定线程模型的参数来调用。CUDA 采用的模型是 2D 网格模型，然后是网格上每个块内部的 3D 线程模型。

对于您遇到的问题类型，除了在每个块中使用 1D 线程的 1D 网格之外，实际上没有必要使用任何东西，因为字符串池像其他问题（例如矩阵乘法）那样拆分成 2D 并没有真正意义

我将通过一个简单的示例，例如字符串池中的 100 个字符串，并且您希望以并行方式而不是按顺序检查它们。

//main
//Should cudamalloc and cudacopy to device up before this code
dim3 dimGrid(10, 1); // 1D grid with 10 blocks
dim3 dimBlocks(10, 1); //1D Blocks with 10 threads 
fun<<<dimGrid, dimBlocks>>>(, Height)
//cudaMemCpy answerIdx back to integer on host

//kernel (Not positive on these types as my CUDA is very rusty
__global__ void fun(char *strings[], char *stringToMatch, int *answerIdx)
{
    int idx = blockIdx.x * 10 + threadIdx.x;

    //Obviously use whatever function you've been using for string comparison
    //I'm just using == for example's sake
    if(strings[idx] == stringToMatch)
    { 
       *answerIdx = idx
    }
}

这显然不是最有效的，并且很可能不是传递参数和使用 CUDA 使用内存的确切方法，但我希望它能够理解拆分工作负载的意义，并且“全局”函数可以在许多不同的地方执行核心，所以你不能真正告诉他们停止。可能有一种我不熟悉的方法，但是通过将工作负载分配到设备上（当然是以合理的方式）所获得的加速已经给你带来了巨大的性能改进。为了了解线程模型，我强烈建议您阅读 Nvidia 网站上的 CUDA 文档。它们将极大地帮助您，并教您设置网格和块以获得最佳性能的最佳方法。

于 2012-09-20T04:02:37.920 回答