在 Cuda 中计算定点逻辑

Question

我是 Cuda 新手，我尝试并行化多个定点操作。 问题是计算固定点的长度可能变化很大，这不是 SIMT 的最佳解决方案。 我尝试将算法切成更小的部分，以便更好地并行化，但我在如何将更小的部分连接在一起的逻辑上遇到了困难。

原始算法的伪代码是这样的：

constant data[big][big];

vec fixpoint(int in) {
    vec v;
    v.push_back(in)
    start = 0
    while start < len(vec) {
        current = v[start]
        for i = 1..n(current, data) {
            if foo(data, v, current)
                v.push_back(data[something])
            if fixpoint_condition(data, v)
                return vec
        }
        start = start + 1
    }
}

int main() {
    for i = 1..M {
        v = fixpoint(i)
        domesomething(v)
    }
}

因此主函数为所有数字 1..M 调用固定点。 这是我可以并行的一点。 但考虑到对固定点的不同调用可能会产生相当大的运行时差异，我认为在扭曲中对固定点进行不同的调用不是一个好主意，因为许多线程不会做一些有用的事情，因为它们已经达到了固定点，而其他线程则无法执行某些有用的操作。线程仍在搜索固定点。 所以我想将我的内核进一步减少到这样的程度：

        current = v[start]
        for i = 1..n(current, data) {
            if foo(data, v, current)
                v.push_back(data[something])
            if fixpoint_condition(data, v)
                return vec
        }

因此每个线程在全局内存中都有一个向量 v （不幸的是我可能必须使用预先分配的静态大小） 并且可以对 v 中的一个元素执行内部 for 循环，从而扩展 v 和/或找到固定点。 现在我可以针对数字 1..M 大规模并行化这些线程。它们都有相似的运行时 (

n(current,data)

)，应该能够最好地利用我的 GPU。 我现在的问题是对外部固定点枚举进行编码。 当线程完成时，CPU（或者可能是 GPU，也许是屏障？）需要检查是否已达到固定点（可能在 v 中设置了标志）或者是否已将新元素添加到 v。如果没有固定点必须安排使用 start+1 到达一个新线程，以继续搜索该特定数字的固定点。

如何使用 CUDA 对该逻辑进行编程。 根据线程的结果，我要么想要安排另一个线程继续处理旧线程的结果。或者停下来并使用结果。

我想过在线程完成后使用cuda evnts与CPU同步。但由于每个线程都需要有不同的事件，我认为这是一个坏主意。另外，cudaEventSynchronize 需要事件作为参数，但我希望回调第一个（以及第二个、第三个等...）完成的任何线程/事件，以便我可以检查是否需要调度新线程。

有人有什么想法吗？我不认为这种问题如此罕见，我猜想在模拟中，也许是光线投射等......你会有具有不同运行时间的固定点操作。我愿意接受有关屏障使用和一般草图的具体建议。

Answer 1

我认为您实际上可以使用 CDP2 从内核中异步启动内核。 从内核调用内核 https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#cuda-dynamic-parallelism

有了这个，实际上应该可以检查结果并动态决定继续内核内的固定点搜索。 （伪代码）

constant data[big][big];

__global__ void subtask(int* data, vec v, int current) {
    for i = 1..n(current, data) {
        if foo(data, v, current)
            v.push_back(data[something])
    }  
}

__global__ void recursion(int* data, vec v, int index) {
    subtask(data, vec, index)
    if !fixpoint_condition(data, v)
        <<<1,1, cudaStreamTailLaunch>>>subtask(data, v, index+1)
    
}

__global__ void fixpoint(int* data) {
    index = threadId.x;
    
    __global__ vec v;
    v.push_back(data[index])
    start = 0
    recursion(data, v, 0)
}

int main() {
        <<<1, M>>>fixpoint(data)
}

定点内核只是开始每个数字 M 的整个进程。 递归函数确实会更改数据，然后决定是否使用 CDP2 和 tailLaunch 继续搜索，并在之后调度另一个任务。这意味着子任务有望在扭曲中对齐。前一个 while 循环中的新迭代确实会获得一个新线程。所有这一切都无需与 CPU 通信。

一个缺点可能是

vec

需要存储在全局内存中，因为不同的线程需要不断访问此内存，并且不能保证这些线程位于一个块中。