CUDA动态并行中的同步

Posted 2021-04-21

我正在使用以下内核测试动态并行性，该内核使用动态并行性以分而治之的方式获取整数数组的最大值：

__global__ void getMax(int * arr, int ini, int fin, int * maxv) {

  if (ini >= fin) return;

  if (fin-ini==1) {

    *maxv = arr[ini];
    printf("Elem: %d (ini:%d)
", *maxv, ini);

  } else {

    int * max1, * max2;
    max1 = (int *) malloc(sizeof(int));
    max2 = (int *) malloc(sizeof(int));

    getMax<<<1,1>>>(arr, ini, (fin+ini)/2, max1);
    getMax<<<1,1>>>(arr, (fin+ini)/2, fin, max2);
    cudaDeviceSynchronize();

    printf("Max1: %d, Max2: %d (ini:%d,fin:%d)
",
        *max1, *max2, ini, fin);
    *maxv = max(*max1, *max2);

    free(max1); free(max2);

  }
}

一个被称为：getMax<<<1,1>>>(d_arr, 0, N, d_max)，d_arr为数组，N为其大小，d_max为最大值。虽然有时我会得到正确的输出，但这个属性我倾向于看错了：

10 6 8 7 14 4 0 4 9 8 6 4 8 10 5 1 
Max1: 0, Max2: 0 (ini:0,fin:4)
Elem: 10 (ini:0)
Max1: 10, Max2: 0 (ini:0,fin:2)
Elem: 6 (ini:1)
Elem: 8 (ini:2)
Max1: 8, Max2: 0 (ini:2,fin:4)
Elem: 7 (ini:3)
Max1: 8, Max2: 8 (ini:4,fin:8)
Elem: 14 (ini:4)
Max1: 14, Max2: 6 (ini:4,fin:6)
Elem: 4 (ini:5)
Elem: 0 (ini:6)
Max1: 0, Max2: 8 (ini:6,fin:8)
Elem: 4 (ini:7)
Max1: 0, Max2: 8 (ini:0,fin:8)
Max1: 0, Max2: 4 (ini:8,fin:12)
Elem: 9 (ini:8)
Max1: 9, Max2: 4 (ini:8,fin:10)
Elem: 8 (ini:9)
Elem: 6 (ini:10)
Max1: 6, Max2: 4 (ini:10,fin:12)
Elem: 4 (ini:11)
Max1: 6, Max2: 6 (ini:12,fin:16)
Elem: 8 (ini:12)
Max1: 8, Max2: 8 (ini:12,fin:14)
Elem: 10 (ini:13)
Elem: 5 (ini:14)
Max1: 5, Max2: 6 (ini:14,fin:16)
Elem: 1 (ini:15)
Max1: 4, Max2: 6 (ini:8,fin:16)
Max1: 8, Max2: 6 (ini:0,fin:16)
Device max: 8
Host max: 14

正如你所看到的，虽然正在使用cudaDeviceSynchronize()，但有很多次父亲网格在他们的孩子完成执行之前打印。更糟糕的是，在最终输出中没有考虑一些子值，从GPU获得错误的结果。

我知道在内核中使用malloc（使用全局内存）和动态并行本身目前还不够快，因为这段代码可以比CPU有更好的加速。我只是想了解为什么这段代码没有正确同步。

答案

每当你在CUDA代码中遇到麻烦时，建议用cuda-memcheck运行你的代码，也可以运行proper CUDA error checking。对于CUDA动态并行（CDP）代码，您可以（并且应该）以相同的方式对设备端内核启动和运行时API的设备使用进行错误检查。即使您不理解生成的错误输出，它也会对那些试图帮助您的人有用。

此外，在寻求无效代码的帮助时，你是supposed提供MCVE。但在这种情况下，我能够通过添加自己的主机测试代码来重新创建您的观察。

在这种情况下的问题似乎是你超过了与CDP相关的默认嵌套和同步深度，正如here所讨论的那样。

通过在主机代码的开头添加一行：

cudaError_t err = cudaDeviceSetLimit(cudaLimitDevRuntimeSyncDepth, 16);

在我的测试用例中，我能够使观察到的问题消失。

但是请注意文档中的最大限制为24，因此在每个递归深度级别需要设备同步的递归机制在问题大小变大时不太可行。

我假设你只是把它作为一种学习练习。如果你真的对有效的最大发现感兴趣，那就有far more efficient reduction based techniques。