__global__ 函数如何像 C/C++ 一样返回值或中断 [英] how can a __global__ function RETURN a value or BREAK out like C/C++ does

问题描述

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

Recently I've been doing string comparing jobs on CUDA, and i wonder how can a __global__ function return a value when it finds the exact string that I'm looking for.

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

I mean, i need the __global__ function which contains a great amount of threads to find a certain string among a big big string-pool simultaneously, and i hope that once the exact string is caught, the __global__ function can stop all the threads and return back to the main function, and tells me "he did it"!

我正在使用 CUDA C.我怎样才能做到这一点?

I'm using CUDA C. How can I possibly achieve this?

推荐答案

在 CUDA(或 NVIDIA GPU)中没有办法让一个线程中断所有正在运行的线程的执行.不能一发现结果就立即退出内核，这在今天是不可能的.

There is no way in CUDA (or on NVIDIA GPUs) for one thread to interrupt execution of all running threads. You can't have immediate exit of the kernel as soon as a result is found, it's just not possible today.

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

But you can have all threads exit as soon as possible after one thread finds a result. Here's a model of how you would do that.

__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 替换为 if 并运行刚好足以覆盖任务数量的线程.那么就没有死锁的机会(但上一点的第一部分适用).
8. workLeftToDo() 是针对问题的，但是当没有工作要做时它会返回false，这样我们就不会在找不到匹配的情况下死锁.

1. Note the use of volatile. This is important.
2. Make sure you initialize found—which must be a device pointer—to false before launching the kernel!
3. Threads will not exit instantly when another thread updates found. They will exit only the next time they return to the top of the while loop.
4. How you implement do_some_work matters. If it is too much work (or too variable), then the delay to exit after a result is found will be long (or variable). If it is too little work, then your threads will be spending most of their time checking found rather than doing useful work.
5. do_some_work is also responsible for allocating tasks (i.e. computing/incrementing indices), and how you do that is problem specific.
6. If the number of blocks you launch is much larger than the maximum occupancy of the kernel on the present GPU, and a match is not found in the first running "wave" of thread blocks, then this kernel (and the one below) can deadlock. If a match is found in the first wave, then later blocks will only run after found == true, which means they will launch, then exit immediately. The solution is to launch only as many blocks as can be resident simultaneously (aka "maximal launch"), and update your task allocation accordingly.
7. If the number of tasks is relatively small, you can replace the while with an if and run just enough threads to cover the number of tasks. Then there is no chance for deadlock (but the first part of the previous point applies).
8. workLeftToDo() is problem-specific, but it would return false when there is no work left to do, so that we don't deadlock in the case that no match is found.

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

Now, the above may result in excessive partition camping (all threads banging on the same memory), especially on older architectures without L1 cache. So you might want to write a slightly more complicated version, using a shared status per block.

__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();
    }
}

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

This way, one thread per block polls the global variable, and only threads that find a match ever write to it, so global memory traffic is minimized.

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

Aside: __global__ functions are void because it's difficult to define how to return values from 1000s of threads into a single CPU thread. It is trivial for the user to contrive a return array in device or zero-copy memory which suits his purpose, but difficult to make a generic mechanism.

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

Disclaimer: Code written in browser, untested, unverified.

这篇关于__global__ 函数如何像 C/C++ 一样返回值或中断的文章就介绍到这了，希望我们推荐的答案对大家有所帮助，也希望大家多多支持IT屋！