CUDA动态并行性中的同步

Question

i正在使用以下内核测试动态并行性，该内核使用动态平行性以分隔和征服方式获得整数数组的最大值：

__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)n", *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)n",
        *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。我只是想理解为什么此代码无法正确同步。

Answer 1

任何时候在CUDA代码中遇到麻烦时，建议使用cuda-memcheck运行代码，并进行适当的CUDA错误检查。对于CUDA动态并行性（CDP）代码，您可以（并且应该）以相同的方式检查运行时API的设备端内核启动和设备使用情况。即使您不了解产生的错误输出，也对那些试图帮助您的人也很有用。

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

在这种情况下，问题似乎是您超出了与CDP相关的默认嵌套和同步深度。

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

cudaError_t err = cudaDeviceSetLimit(cudaLimitDevRuntimeSyncDepth, 16);

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

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

我认为您只是将其作为学习练习。如果您实际上对有效的最大法合感兴趣，则基于减少的技术有更高的效率。