CUDA：如何在指向数组的指针数组上应用__restrict __？

Question

此内核使用两个__restrict__ int数组可以正常编译：

__global__ void kerFoo( int* __restrict__ arr0, int* __restrict__ arr1, int num )
{
    for ( /* Iterate over array */ )
        arr1[i] = arr0[i];  // Copy one to other
}

但是，组成指针数组的相同两个int数组编译失败：

__global__ void kerFoo( int* __restrict__ arr[2], int num )
{
    for ( /* Iterate over array */ )
        arr[1][i] = arr[0][i];  // Copy one to other
}

编译器给出的错误是：

error: invalid use of `restrict'

我有某些结构，这些结构由指向数组的指针组成。 （例如，传递给具有int* arr[16]的内核的结构。）如何将它们传递给内核并能够在其上应用__restrict__？

Answer 1

通过一些任意迭代填充代码中的注释，我们得到以下程序：

__global__ void kerFoo( int* __restrict__ arr[2], int num )
{
    for ( int i = 0; i < 1024; i ++)
        arr[1][i] = arr[0][i];  // Copy one to other
}

以及带有CUDA 10.1的compiles fine（Godbolt.org）。

Answer 2

[CUDA C手册仅引用[99]的C99定义，没有特殊的CUDA特定情况。

由于指示的参数是一个包含两个指针的数组，因此对我来说，对__restrict__的这种用法看起来非常有效，编译器没有理由抱怨恕我直言。我将要求编译器作者验证并可能/可能更正此问题。不过，我会对不同的观点感兴趣。

对@talonmies的一句话：

restrict的全部要点是告诉编译器两个或多个指针参数在内存中永远不会重叠。

严格来说这不是真的。 __restrict__告诉编译器，所讨论的指针在其生命周期内是可以访问指向对象的唯一指针。请注意，指向的对象只是assumed是restrict的数组。 （实际上，在这种情况下，它只有一个int。）由于编译器无法知道数组的大小，因此程序员必须保护数组的边界。