CUDA 动态并行中的同步

Question

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

__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 运行您的代码并执行 proper CUDA error checking。对于 CUDA 动态并行 (CDP) 代码，您可以（并且应该）以相同的方式对运行时间 API 的设备端内核启动和设备使用情况进行错误检查。即使您不理解生成的错误输出，它也会对那些试图帮助您的人有用。

此外，当您就无法使用的代码寻求帮助时，您是 supposed to provide a MCVE。但在这种情况下，我能够通过添加我自己的主机测试代码来重新创建您的观察结果。

本例中的问题似乎是您超出了与 CDP 关联的默认嵌套和同步深度，如 here 所述。

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

cudaError_t err = cudaDeviceSetLimit(cudaLimitDevRuntimeSyncDepth, 16);

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

但是请注意文档中的最大限制是 24，因此当您的问题规模变大时，您的递归机制要求在每个递归深度级别同步设备并不是很可行。

我假设您只是将此作为一种学习练习。如果你真的对高效的求最大值感兴趣，有 far more efficient reduction based techniques.

CUDA 动态并行中的同步

Synchronization in CUDA dynamic parallelism

cuda

dynamic-parallelism