为什么我收到错误消息：不允许“限制”？答案

【问题标题】：Why am I getting the error message: "restrict" not allowed?为什么我收到错误消息：不允许“限制”？
【发布时间】：2013-12-04 14:54:21
【问题描述】：

我正在编写一个 CUDA 内核并想__restrict__ 我的一些参数。我收到错误消息：

"restrict" is not allowed

某些变量类型不允许吗？对于一些参数组合？因为一些编译器标志？因为我一直很淘气？

简化的内核签名：

template <typename T> foo(
    const T a[],
    __restrict__ SomeType b[],
    const T c
) {
    /* etc. */
}

【问题讨论】：

【解决方案1】：

您只能在指针类型上使用__restrict__。这也是 __restrict__ 有意义的唯一上下文。

【讨论】：

只是为了放大：__restrict__ SomeType b[] 不行。 SomeType __restrict__ b[] 不行。 __restrict__ SomeType *b 不行。 SomeType* __restrict__ b 没问题。参考documentation的例子。
那么，你是说如果我想要一个指向 SomeType 的“受限”指针，我就不能对参数使用方括号表示法吗？
@einpoklum：您不能修改数组引用的值（我不确定这是否是正确的术语）。也就是说，您不能对b 执行任何指针运算，并且它将始终指向内存中的同一点。这就是为什么别名问题不适用于b。您可以通过获取b 的地址来制作指针。您可以在该新指针上使用__restrict__。
@einpoklum：我刚刚意识到，如果您在同一个数组中使用多个索引，您可能会遇到别名问题，所以我现在不确定为什么 __restrict__ 不能应用于数组。也许编译器已经假定数组没有别名。必须研究...
formal definition 的 restrict 在 c99 标准中给出（特别注意链接的第 3.10 节）。 restrict 是一个指针限定符。它不限定指向的对象。因此，在遇到restrict 的那一刻，我们一定指的是一个指针类型。 cuda __restrict__ 限定符模仿了这种行为。所以如果T 不是指针类型，那么T restrict 将不起作用，不管它后面是什么。你也不能按照你在 C 中的建议去做。这不是真正的 CUDA 问题。

【解决方案2】：

如果您在指向函数的指针参数上使用了__restrict__，但顺序错误，也会发生这种情况，例如：

__device__ int bar(float * __restrict__ x);

【讨论】：