1

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

内核需要一些临时数据和共享内存，但由于它很大，它使用 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();
}

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

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

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

我对动态并行性不太怀疑，因为我可以轻松地保持大循环运行，但我不确定我是否可以在这里拥有优势。

4

1 回答 1

2

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

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)内核中。

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

于 2013-06-17T13:36:46.567 回答