我是 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 需要事件作为参数,但我希望回调第一个(以及第二个、第三个等...)完成的任何线程/事件,以便我可以检查是否需要调度新线程。
有人有什么想法吗?我不认为这种问题如此罕见,我猜想在模拟中,也许是光线投射等......你会有具有不同运行时间的固定点操作。我愿意接受有关屏障使用和一般草图的具体建议。
我认为您实际上可以使用 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
需要存储在全局内存中,因为不同的线程需要不断访问此内存,并且不能保证这些线程位于一个块中。