在 CUDA（或 NVIDIA GPU）中，一个线程无法中断所有正在运行的线程的执行。你不能在发现结果后立即退出内核，这在今天是不可能的。
但你可以让所有线程退出尽快地当一个线程找到结果后。这是一个如何做到这一点的模型。
__global___ void kernel(volatile bool *found, ...)
{
while (!(*found) && workLeftToDo()) {
bool iFoundIt = do_some_work(...); // see notes below
if (iFoundIt) *found = true;
}
}
关于此的一些注释。
- 注意使用
volatile。这个很重要。
- 确保初始化
found— 必须是一个设备指针 — 到false在启动内核之前！
- 当另一个线程更新时，线程不会立即退出
found。仅当它们下次返回到 while 循环顶部时才会退出。
- 你如何实施
do_some_work很重要。如果工作量太大（或变化太大），那么找到结果后退出的延迟将会很长（或变化）。如果工作量太少，那么您的线程将花费大部分时间进行检查found而不是做有用的工作。
-
do_some_work还负责分配任务（即计算/递增索引），以及如何做到这一点是特定于问题的。
- 如果您启动的块数远大于当前 GPU 上内核的最大占用率，并且在第一个运行的线程块“波”中未找到匹配项，则该内核（以及下面的内核）可能会死锁。如果在第一波中找到匹配项，则后面的块将仅在之后运行
found == true，这意味着它们将启动，然后立即退出。解决方案是仅启动可同时驻留的尽可能多的块（也称为“最大启动”），并相应地更新任务分配。
- 如果任务数量比较少，可以更换
while与if并运行足够的线程来覆盖任务数量。这样就不会出现死锁（但上一点的第一部分适用）。
-
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__ 函数是无效的，因为很难定义如何将数千个线程的值返回到单个 CPU 线程中。对于用户来说，在设备或零拷贝内存中设计一个适合其目的的返回数组很简单，但很难建立通用机制。
免责声明：在浏览器中编写的代码，未经测试，未经验证。
