Why am I getting the error message: "restrict" not allowed?

Question

I'm writing a CUDA kernel and want to __restrict__ some of my parameters. I'm getting the error message:

"restrict" is not allowed

Is it not allowed for some variable types? For some combinations of parameters? Because of some compiler flags? Because I've been naughty?

Simplified kernel signature:

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

score 7 · Accepted Answer · answered Dec 04 '13 at 16:59

7

You can only use __restrict__ on a pointer type. That is also the only context in which __restrict__ makes sense.

answered Dec 04 '13 at 16:59

Roger Dahl

15,132
8
62
82

14

Just to amplify: `__restrict__ SomeType b[]` is not OK. `SomeType __restrict__ b[]` is not OK. `__restrict__ SomeType *b` is not OK. `SomeType* __restrict__ b` is OK. Refer to the [documentation](http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#restrict) for an example. – Robert Crovella Dec 04 '13 at 17:35
So, you're saying that if I want a 'restrict'ed pointer to SomeType, there's no way I can use square bracket notation for the parameter? – einpoklum Dec 05 '13 at 09:28
@einpoklum: You can't modify the value of the array reference (I'm not sure if that is the correct terminology). That is, you can't perform any pointer arithmetic on `b` and it will always point to the same point in memory. Which is why the aliasing issue does not apply to `b`. You can make a pointer by taking the address of `b`. You can use `__restrict__` on that new pointer. – Roger Dahl Dec 05 '13 at 14:36
@einpoklum: I just realized that you can get aliasing issues if you use multiple indexes into the same array, so I'm not sure now why `__restrict__` can't be applied to an array. Maybe the compiler already assumes no aliasing for arrays. Would have to research... – Roger Dahl Dec 05 '13 at 15:02
I know that in C, `void foo(int a[])` is define to be exactly `void foo(int* a)`; that is, you can't pass arrays anyway, only pointers. Isn't that also the case in C++? – einpoklum Dec 05 '13 at 15:32
@einpoklum: Yes, it is the case -- you can't pass arrays in C++ either. Arrays are implicitly converted to pointers. But arrays and pointers have different semantics. See for instance, http://c-faq.com/aryptr/index.html and http://www.tutorialspoint.com/cplusplus/cpp_pointers_vs_arrays.htm. – Roger Dahl Dec 05 '13 at 18:19
I know very well about the differences. The point is, if you can't pass arrays, how can `T* __restrict__ ptr` be Ok while `T __restrict__ ptr[]` not be Ok? – einpoklum Dec 05 '13 at 18:21
@einpoklum: I fail to see how those two are related...? Let's move to chat if you want to clarify. But, as I mentioned, I can see a use case for `__restrict__` on arrays unless the compiler assumes that there is no aliasing for arrays. Beyond that, I would have to research. – Roger Dahl Dec 05 '13 at 20:28
3

A [formal definition](http://www.lysator.liu.se/c/restrict.html#formal-definition) of `restrict` is given in the c99 standard (especially note the linked section 3.10). `restrict` is a *pointer qualifier*. It does not qualify the object pointed to. Therefore, at the moment `restrict` is encountered, we must be referring to a *pointer type*. The cuda `__restrict__` qualifier mimics this behavior. So if `T` is not a pointer type, then `T restrict` will not work, regardless of what follows it. You cannot do what you are suggesting in C, either. This is not really a CUDA issue. – Robert Crovella Dec 05 '13 at 20:54
@RobertCrovella: Got it. Ok then. – einpoklum Dec 05 '13 at 21:31

einpoklum · Answer 2 · 2023-08-17T21:47:31.533

3

This can also happen if you used __restrict__ on a pointer parameter to a function, but got the order wrong, e.g.:

__device__ int bar(float __restrict__ *  x); // wrong order!

instead of:

__device__ int bar(float * __restrict__ x); // correct order

edited Aug 17 '23 at 21:47

answered Mar 07 '21 at 17:12

einpoklum

118,144
57
340
684

Please explain a bit more in detail: Did you show a correct order or an example of a wrong one? Can you write one example of a correct one and one example of a wrong one and explain why the wrong one is wrong? Thank you! – ElectRocnic Aug 17 '23 at 11:22
1

@ElectRocnic: You're quite right, see edit. – einpoklum Aug 17 '23 at 21:47

Why am I getting the error message: "restrict" not allowed?

2 Answers2