Is the restriction (amp) more stringent than the CUDA kernel code?

Question

In C ++, AMP kernel functions or lambdas are marked with a restriction (amp), which imposes severe restrictions on the permitted subset of C ++ ( listed here ). Does CUDA provide more freedom for a subset of C or C ++ in kernel functions?

Answer 1

As with Visual Studio 11 and CUDA 4.1, restrict(amp) functions are more stringent than similar CUDA __device__ functions. Most notably, AMP is more restrictive of how pointers can be used. This is a natural consequence of the computing substrate AMP DirectX11, which prohibits pointers in HLSL (graphical shader). In contrast, the lower-level IRU is CUDA PTX , which is a more general goal than HLSL.

Here's a string-by-string comparison:

 | VS 11 AMP restrict(amp) functions | CUDA 4.1 sm_2x __device__ functions | |------------------------------------------------------------------------------| |* can only call functions that have |* can only call functions that have | | the restrict(amp) clause | the __device__ decoration | |* The function must be inlinable |* need not be inlined | |* The function can declare only |* Class types are allowed | | POD variables | | |* Lambda functions cannot |* Lambdas are not supported, but | | capture by reference and | user functors can hold pointers | | cannot capture pointers | | |* References and single-indirection |* References and multiple-indirection | | pointers are supported only as | pointers are supported | | local variables and function | | |* No recursion |* Recursion OK | |* No volatile variables |* Volatile variables OK | |* No virtual functions |* Virtual functions OK | |* No pointers to functions |* Pointers to functions OK | |* No pointers to member functions |* Pointers to member functions OK | |* No pointers in structures |* Pointers in structures OK | |* No pointers to pointers |* Pointers to pointers OK | |* No goto statements |* goto statements OK | |* No labeled statements |* Labeled statements OK | |* No try, catch, or throw statements |* No try, catch, or throw statements | |* No global variables |* Global __device__ variables OK | |* Static variables through tile_static |* Static variables through __shared__ | |* No dynamic_cast |* No dynamic_cast | |* No typeid operator |* No typeid operator | |* No asm declarations |* asm declarations (inline PTX) OK | |* No varargs |* No varargs | 

Read more about restrict(amp) restrictions here . You can read about C ++ support in the CUDA __device__ functions in Appendix D of the CUDA C Programming Guide .