0

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

__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 上有很好的加速。我只是想了解为什么这段代码没有正确同步。

4

1 回答 1

4

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

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

这种情况下的问题似乎是您超出了与 CDP 关联的默认嵌套和同步深度，如此处所述。

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

cudaError_t err = cudaDeviceSetLimit(cudaLimitDevRuntimeSyncDepth, 16);

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

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

我假设您只是将其作为学习练习。如果您真的对有效的最大值查找感兴趣，那么还有更有效的基于归约的技术。

于 2017-12-23T23:10:52.553 回答