__global__函数如何能像C/C++那样返回一个值或BREAK

Question

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

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

我正在使用CUDA C.我怎么可能实现这个目标？

Answer 1

在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上内核的最大占用率,并且在第一个运行的"wave"线程块中找不到匹配,则此内核(以及下面的内核)可能会死锁.如果在第一个波浪中找到匹配,则后面的块将仅在之后运行found == true,这意味着它们将启动,然后立即退出.解决方案是仅启动可以同时驻留的块(也称为"最大启动"),并相应地更新任务分配.
7. 如果任务的数量比较少,可以代替while用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线程中.用户设计适合其目的的设备或零拷贝存储器中的返回数组是微不足道的,但很难制作通用机制.

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

Answer 2

如果您有冒险精神,停止内核执行的另一种方法就是执行

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

如果找到答案.

这不需要轮询内存,但是不如Mark Harris建议的解决方案,因为它使内核以错误条件退出.这可能会掩盖实际错误(因此请务必以明确允许从错误中成功执行的方式写出结果),并且可能会导致其他打嗝或降低整体性能,因为驱动程序会将此视为异常.

如果您寻求安全而简单的解决方案,请转而选择Mark Harris的建议.