GPU上的独立搜索——如何同步完成？

Question

假设我有一些算法 generateRandomNumbersAndTestThem()，它以概率 p 返回 true，以概率 1-p 返回 false。通常 p 非常小，例如p=0.000001。

我正在尝试在 JOCL 中构建一个程序，该程序估计 p 如下： generateRandomNumbersAndTestThem() 在所有可用的着色器核心（最好是多个 GPU）上并行执行，直到找到至少 100 个真值。那么 p 的估计值为 100/n，其中 n 是 generateRandomNumbersAndTestThem() 被执行的总次数。

对于 p = 0.0000001，这意味着大约 10^9 次独立尝试，这应该很明显为什么我希望在 GPU 上执行此操作。但我正在努力如何正确实施停止条件。我的想法是有一些类似的东西作为内核：

__kernel void sampleKernel(all_the_input, __global unsigned long *totAttempts) {
    int gid = get_global_id(0);
    //here code that localizes all_the_input for faster access
    while (lessThan100truesFound) {
        totAttempts[gid]++;
        if (generateRandomNumbersAndTestThem()) 
            reportTrue();
    }
}

考虑到

，我应该如何在没有严重性能损失的情况下实现这一点

• “if”的触发将是一个非常罕见的事件，因此如果所有线程都必须在执行 reportTrue() 时等待，这不是问题
• 当第 100 次调用 reportTrue() 时，lessThan100truesFound 只需修改一次（从 true 到 false）（所以我什至不知道布尔值是否正确）
• 计划为此购买全新的 GPU 硬件，因此您可以假设最近的 GPU，例如多个 ATI Radeon HD7970。但如果我能在我目前的 HD5450 上进行测试，那就太好了。

我假设可以完成类似于 Java 的“同步”修饰符的操作，但我找不到确切的方法。执行此操作的“正确”方法是什么，即在没有严重性能损失的情况下工作的任何方法？

Answer 1

我建议不要使用全局标志来停止内核，而是运行内核进行一定数量的尝试，检查主机是否积累了足够的“成功”，并在必要时重复。在内核中使用未定义长度的循环是不好的，因为 GPU 驱动程序可能会被看门狗定时器杀死。此外，在每次迭代时检查一些全局变量肯定会影响内核性能。

这样，reportTrue 可以作为atomic_inc 实现到驻留在全局内存中的某个计数器。

__kernel void sampleKernel(all_the_input, __global unsigned long *successes) {
    int gid = get_global_id(0);
    //here code that localizes all_the_input for faster access
    for (int i = 0; i < ATT_PER_THREAD; ++i) {
        if (generateRandomNumbersAndTestThem()) 
            atomic_inc(successes);
    }
}

ATT_PER_THREAD 将根据执行generateRandomNumbersAndTestThem() 所需的时间进行调整。内核启动开销非常小，因此通常不需要让您的内核运行超过 0.1--1 秒