@@ -33,12 +33,14 @@ The following Python constructs are not supported:
3333* Comprehensions (either list, dict, set or generator comprehensions)
3434* Generator (any ``yield `` statements)
3535
36- The ``raise `` statement is supported.
36+ The ``raise `` and ``assert `` statements are supported, with the following
37+ constraints:
3738
38- The ``assert `` statement is supported, but only has an effect when
39- ``debug=True `` is passed to the :func: `numba.cuda.jit ` decorator. This is
40- similar to the behavior of the ``assert `` keyword in CUDA C/C++, which is
41- ignored unless compiling with device debug turned on.
39+ - They can only be used in kernels, not in device functions.
40+ - They only have an effect when ``debug=True `` is passed to the
41+ :func: `@cuda.jit <numba.cuda.jit> ` decorator. This is similar to the behavior
42+ of the ``assert `` keyword in CUDA C/C++, which is ignored unless compiling
43+ with device debug turned on.
4244
4345
4446Printing of strings, integers, and floats is supported, but printing is an
@@ -54,6 +56,56 @@ This is due to a general limitation in CUDA printing, as outlined in the
5456<https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#limitations> `_
5557in the CUDA C++ Programming Guide.
5658
59+
60+ Recursion
61+ ---------
62+
63+ Self-recursive device functions are supported, with the constraint that
64+ recursive calls must have the same argument types as the initial call to
65+ the function. For example, the following form of recursion is supported:
66+
67+ .. code :: python
68+
69+ @cuda.jit (" int64(int64)" device = True )
70+ def fib (n ):
71+ if n < 2 :
72+ return n
73+ return fib(n - 1 ) + fib(n - 2 )
74+
75+ fib `` function always has an ``int64 `` argument), whereas the following
76+ is unsupported:
77+
78+ .. code :: python
79+
80+ # Called with x := int64, y := float64
81+ @cuda.jit
82+ def type_change_self (x , y ):
83+ if x > 1 and y > 0 :
84+ return x + type_change_self(x - y, y)
85+ else :
86+ return y
87+
88+ type_change_self `` provides ``(int64, float64) `` arguments,
89+ but the inner call uses ``(float64, float64) `` arguments (because ``x - y `` /
90+ ``int64 - float64 `` results in a ``float64 `` type). Therefore, this function is
91+ unsupported.
92+
93+ Mutual recursion between functions (e.g. where a function ``func1() `` calls
94+ ``func2() `` which again calls ``func1() ``) is unsupported.
95+
96+ .. note ::
97+
98+ The call stack in CUDA is typically quite limited in size, so it is easier
99+ to overflow it with recursive calls on CUDA devices than it is on CPUs.
100+
101+ Stack overflow will result in an Unspecified Launch Failure (ULF) during
102+ kernel execution. In order to identify whether a ULF is due to stack
103+ overflow, programs can be run under `Compute Sanitizer
104+ <https://docs.nvidia.com/compute-sanitizer/ComputeSanitizer/index.html> `_,
105+ which explicitly states when stack overflow has occurred.
106+
107+ .. _cuda-built-in-types :
108+
57109Built-in types
58110===============
59111
0 commit comments