0

终于启动并运行了动态并行，我现在正尝试用它来实现我的模型。我花了一段时间才发现需要使用 cudaDeviceSynchronize() 使父内核等待子内核完成导致了一些奇怪的输出。

我定义为 arrAdd 的设备函数似乎有问题。这是 k2 父内核中每个子内核之前和之后的输出表。

Initially    : k1   = { -1   0   0   0   0 }
Post arrInit : temp = { .25 .25 .25 .25 .25}
Post arrMult : temp = {-.25  0   0   0   0 }
post arrAdd  : temp = { -8   0   0   0   0 }
Expected     : temp = {-.50  0   0   0   0 }


__global__ void k2(double* concs, int* maxlength, double* k1s, double* k2s, double * temp, double* tempsum)
{
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    double a21 = .25;

    arrInit<<< 1, *maxlength >>>(temp, a21);                //temp = a21
    cudaDeviceSynchronize();
    arrMult<<< 1, *maxlength >>>(k1s, temp, temp);          //temp = a21*k1
    cudaDeviceSynchronize();
    arrAdd<<< 1, *maxlength >>>(temp, temp, temp);          //temp = 2*a21*k1
    cudaDeviceSynchronize();
}

__global__ void arrAdd(double* a, double* b, double* c)
{
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    c[idx]=a[idx]+b[idx];
}
__global__ void arrMult(double* a, double* b, double* c)
{
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    c[idx]=a[idx]*b[idx];
}
__global__ void arrInit(double* a, double b)
{
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    a[idx]=b;
}
4

1 回答 1

2

您可能不需要与父内核同步。子内核按照父内核指定的顺序执行，父内核的结尾是与最后一个子内核的隐式同步点。

使用动态并行时，请注意以下事项：

  1. 你可以去的最深是 24 (CC=3.5)。

  2. 同时等待启动的动态内核的数量是有限的（CC=3.5 时默认为 2048），但可以增加。

  3. 在子内核调用之后让父内核保持忙碌，否则很可能会浪费资源。

我猜你奇怪的错误结果源于上面提到的第二个因素。当您达到限制时，一些动态内核根本不会运行，如果您不检查错误，您将不会注意到，因为错误创建机制是每个线程的。

您可以通过以cudaLimitDevRuntimePendingLaunchCount作为限制的cudaDeviceSetLimit()来增加此限制。但是你指定的越多，你消耗的全局内存空间就越多。在此处查看文档的 C.4.3.1.3 部分。

于 2013-10-22T21:23:39.210 回答