CUDA：如何在指向数组的指针数组上应用 __restrict__？答案

【问题标题】：CUDA: How to apply __restrict__ on array of pointers to arrays?CUDA：如何在指向数组的指针数组上应用 __restrict__？
【发布时间】：2011-12-07 07:35:06
【问题描述】：

这个使用两个__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__？

【问题讨论】：

__restrict__ 对于指针确实很有用，但 int* arr[2] 实际上是一个数组，如果有两个点。我认为，它不适用于主机代码...
Restrict 真的能给你带来一些性能上的好处吗？？
第二次使用__restrict__ 完全没有意义。 __restrict__ 的全部意义在于告诉编译器两个或多个指针参数永远不会在内存中重叠。在这种情况下，您没有两个指针参数，因此 __restrict__ 不适用。
为什么不int * __restrict__ * __restrict__ arr？不过，这可能太过分了。你可能只需要int * __restrict__ * arr。

标签： cuda restrict-qualifier

【解决方案1】：

CUDA C 手册仅参考了__restrict__ 的C99 定义，没有针对CUDA 的特殊情况。

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

对@talonmies 的评论：

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

这并不完全正确。 restrict 告诉编译器，在其生命周期内，所讨论的指针是唯一可以访问指向对象的指针。请注意，指向的对象仅假定为int 的数组。（实际上，在这种情况下它只有一个int。）由于编译器无法知道数组的大小，因此由程序员来保护数组的边界..

【讨论】：

【解决方案2】：

用一些任意迭代在你的代码中填写注释，我们得到以下程序：

__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
}

还有这个 compiles fine 与 CUDA 10.1 (Godbolt.org)。

【讨论】：