拆分大 CUDA 内核和使用动态并行的好处

Question

我有一个大内核，其中初始状态是使用不同的技术演变而来的。也就是说，我在内核中有一个循环，在这个循环中，某个谓词会根据当前状态进行评估，并根据该谓词的结果执行某个操作。

内核需要一些临时数据和共享内存，但由于它很大，它使用 63 个寄存器，占用率非常低。

我想将内核拆分为许多小内核，但每个块都完全独立于其他块，我（想我）不能在主机代码上使用单个线程来启动多个小内核。

我不确定流是否足以胜任此类工作，我从未使用过它们，但由于我可以选择使用动态并行性，我希望这是实现此类工作的好选择。 从内核启动内核是否很快？ 我是否需要复制全局内存中的数据以供子内核使用？

如果我将我的大内核分成许多小内核，并让第一个内核保留一个主循环，在必要时调用所需的内核（这允许我在每个子内核中移动临时变量），将帮助我增加入住率？

我知道这是一个有点笼统的问题，但我不知道这项技术，我想知道它是否适合我的情况或者流是否更好。

编辑： 为了提供一些其他细节，您可以想象我的内核具有这种结构：

__global__ void kernel(int *sampleData, int *initialData) {
    __shared__ int systemState[N];
    __shared__ int someTemp[N * 3];
    __shared__ int time;
    int tid = ...;
    systemState[tid] = initialData[tid];

    while (time < TIME_END) {
        bool c = calc_something(systemState);
        if (c)
            break;
        someTemp[tid] = do_something(systemState);
        c = do_check(someTemp);
        if (__syncthreads_or(c))
            break;
        sample(sampleData, systemState);
        if (__syncthreads_and(...)) {
            do_something(systemState);
            sync();
            time += some_increment(systemState);
        }
        else {
            calcNewTemp(someTemp, systemState);
            sync();
            do_something_else(someTemp, systemState);
            time += some_other_increment(someTemp, systemState);
        }
    }
    do_some_stats();
}

这是为了告诉你有一个主循环，有临时数据在某处使用而不是在其他点，有共享数据、同步点等。

线程用于计算矢量数据，而理想情况下，每个块中都有一个循环（当然，这不是真的，但从逻辑上讲是这样）......每个块都有一个“大流”。

现在，我不确定在这种情况下如何使用流...“大循环”在哪里？我猜在主机上......但是我如何从一个循环中协调所有块？这是最让我怀疑的地方。我可以使用来自不同主机线程的流吗（每个块一个线程）？

我对动态并行性不太怀疑，因为我可以很容易地保持大循环运行，但我不确定我是否可以在这里获得优势。

Answer 1

我受益于动态并行性来解决以下形式的插值问题：

int i = threadIdx.x + blockDim.x * blockIdx.x;

for(int m=0; m<(2*K+1); m++) {

    PP1 = calculate_PP1(i,m);
    phi_cap1 = calculate_phi_cap1(i,m);  

        for(int n=0; n<(2*K+1); n++) {

            PP2 = calculate_PP2(i,m);
            phi_cap2 = calculate_phi_cap2(i,n);

            atomicAdd(&result[PP1][PP2],data[i]*phi_cap1*phi_cap2); } } }

在哪里K=6。在这个插值问题中，每个加数的计算都是独立于其他的，所以我将它们拆分到一个(2K+1)x(2K+1) 内核中。

根据我（可能不完整）的经验，如果您有几个独立的迭代，动态并行会有所帮助。对于大量的迭代，也许您最终可能会多次调用子内核，因此您应该检查内核启动的开销是否会成为限制因素。