拆分大型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)核。

从我的(可能不完整的)经验来看，如果你有一些独立的迭代，动态并行将会有所帮助。对于大量的迭代，您可能需要多次调用子内核，因此您应该检查内核启动的开销是否会成为限制因素。