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

CUDA C 手册只引用了 __restrict__ 的 C99 定义，没有特殊的 CUDA 特定情况。

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

对@talonmies 的评论:

The whole point of restrict is to tell the compiler that two or more pointer arguments will never overlap in memory.

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