Repository: numba/numba
Pull request: Fall back on inspect.getsourcelines() in first line finder
This enables lineinfo to be produced for cells in Jupyter notebooks.

This was initially manually verified with a kernel in source in a .py file and the same in a Jupyter notebook, with `lineinfo=True` - these two cases produce PTX with identical lineinfo (`.loc` directives).

A test that modifies the linecache in a similar manner to the Jupyter notebook is added. This is based on the `test_string_eval` test, but the difference is that the first line finder does succeed in finding the first line from the linecache.

Fixes #9385.

Problem statement:
Cannot find source for CUDA kernels decorated with `lineinfo=True` when executed from Jupyter notebook cell
For CUDA profiling and potentially debugging purposes, it is important to be able to correlate high-level numba Python code with low-level assembly. This is achieved by providing lineinfo during compilation. In CUDA kernels written in numba cuda, functions have to be decorated with `lineinfo=True` for this.

When executing such functions in standalone Python scripts, profiling tools like Nsight Compute are provided this lineinfo. However, when executing such functions from a cell in a jupyter notebook, numba issues a warning that no function source can be found. This prevents tools integrated with jupyter to fully profile such kernels, since this correlation from SASS assembly to high-level Python is not possible anymore.

```
import math
import numba.cuda
import numpy as np

(width, height) = (128, 128)
np_img = np.zeros((width, height, 1), dtype=int)
gpu_mem = numba.cuda.to_device(np_img)

@numba.cuda.jit(lineinfo=True)
def kernel(buf):
    tx, ty = numba.cuda.grid(2)
    buf[ty][tx] = 1

block = (8,8)    
grid = (math.ceil(width/block[0]), math.ceil(height/block[1]))

kernel[grid, block](gpu_mem)
numba.cuda.synchronize()
```
```
/usr/local/lib/python3.8/dist-packages/numba/core/lowering.py:112: NumbaDebugInfoWarning: Could not find source for function: <function kernel at 0x7f6e5a812700>. Debug line information may be inaccurate.
  warnings.warn(NumbaDebugInfoWarning(msg))
  ```
This was tested with numba 0.58 and Jupyterlab 4.0.10.

Hints from issue discussion:
We need to rework the way we get source line info to use `inspect.getsourcelines()` (as the error reporting and caching mechanisms use) to be able to get the source line info - the present implementation attempts to read the line info from the file, which doesn't exist for Jupyter notebook cells.

Marking as a feature request, with the aim of resolving it for the release following 0.59 (which is already in RC).
This issue is marked as stale as it has had no activity in the past 30 days. Please close this issue if no further response or action is needed. Otherwise, please respond with any updates and confirm that this issue still needs to be addressed.
Still required to support tools working on jupyter cells directly.

Commit message: First line finder: be more explicit about returning None

The first line finder returns *something*, and if we didn't find the
first line then that thing is `None`. This change makes it more clear
that we're doing that, rather than just aborting our attempt early with
a plain `return`.

Patch to review:
diff --git a/docs/upcoming_changes/9861.improvement.rst b/docs/upcoming_changes/9861.improvement.rst
new file mode 100644
index 00000000000..3212a9b5b58
--- /dev/null
+++ b/docs/upcoming_changes/9861.improvement.rst
@@ -0,0 +1,7 @@
+Fall back to ``inspect.getsourcelines()`` in first line finder
+--------------------------------------------------------------
+
+If the first line finder can't locate source code on disk, it now tries using
+``inspect.getsourcelines()`` to get the source. This enables it to locate
+source for cells in Jupyter notebooks, so that line info for them can be
+generated for profiling tools.
diff --git a/numba/misc/firstlinefinder.py b/numba/misc/firstlinefinder.py
index 83558d58ee1..07950bafd48 100644
--- a/numba/misc/firstlinefinder.py
+++ b/numba/misc/firstlinefinder.py
@@ -4,6 +4,7 @@
 """
 
 import ast
+import inspect
 
 
 class FindDefFirstLine(ast.NodeVisitor):
@@ -15,15 +16,15 @@ class FindDefFirstLine(ast.NodeVisitor):
         Or, ``None`` if the definition is not found.
     """
 
-    def __init__(self, code):
+    def __init__(self, name, firstlineno):
         """
         Parameters
         ----------
         code :
             The function's code object.
         """
-        self._co_name = code.co_name
-        self._co_firstlineno = code.co_firstlineno
+        self._co_name = name
+        self._co_firstlineno = firstlineno
         self.first_stmt_line = None
 
     def _visit_children(self, node):
@@ -82,11 +83,21 @@ def get_func_body_first_lineno(pyfunc):
     co = pyfunc.__code__
     try:
         with open(co.co_filename) as fin:
-            file_content = fin.read()
+            source = fin.read()
+            offset = 0
     except (FileNotFoundError, OSError):
-        return
+        try:
+            lines, offset = inspect.getsourcelines(pyfunc)
+            source = "".join(lines)
+            offset = offset - 1
+        except (OSError, TypeError):
+            return None
+
+    tree = ast.parse(source)
+    finder = FindDefFirstLine(co.co_name, co.co_firstlineno - offset)
+    finder.visit(tree)
+    if finder.first_stmt_line:
+        return finder.first_stmt_line + offset
     else:
-        tree = ast.parse(file_content)
-        finder = FindDefFirstLine(co)
-        finder.visit(tree)
-        return finder.first_stmt_line
+        # No first line found.
+        return None
diff --git a/numba/tests/test_firstlinefinder.py b/numba/tests/test_firstlinefinder.py
index 497f148ed06..0608740d0c3 100644
--- a/numba/tests/test_firstlinefinder.py
+++ b/numba/tests/test_firstlinefinder.py
@@ -1,4 +1,5 @@
 import unittest
+import linecache
 import inspect
 
 from numba import njit
@@ -79,6 +80,32 @@ def test_string_eval(self):
         # Cannot determine first line of string evaled functions
         self.assertIsNone(first_def_line)
 
+    def test_string_eval_linecache(self):
+        source = """def foo():
+            pass
+        """
+
+        # Modify the line cache in a similar manner to Jupyter, so that
+        # get_func_body_first_lineno can find the code using the fallback to
+        # inspect.getsourcelines()
+        filename = "<foo>"
+        timestamp = None
+        entry = (len(source), timestamp, source.splitlines(), filename)
+        linecache.cache[filename] = entry
+
+        # We need to compile the code so we can give it the fake filename used
+        # in the linecache
+        code = compile(source, filename, "exec")
+
+        globalns = {}
+        exec(code, globalns)
+        foo = globalns['foo']
+
+        # We should be able to determine the first line number even though the
+        # source does not exist on disk
+        first_def_line = get_func_body_first_lineno(foo)
+        self.assertEqual(first_def_line, 1)
+
     def test_single_line_function(self):
         @njit
         def foo(): pass   # noqa: E704


Code review comments:
diff_hunk: @@ -79,6 +80,32 @@ def test_string_eval(self):
         # Cannot determine first line of string evaled functions
         self.assertIsNone(first_def_line)
 
+    def test_string_eval_linecache(self):
+        source = """def foo():
+            pass
+        """
+
+        # Modify the line cache in a similar manner to Jupyter, so that
+        # get_func_body_first_lineno can find the code using the fallback to
+        # inspect.getsourcelines()
+        filename = "<foo>"
+        timestamp = None
+        entry = (len(source), timestamp, source.splitlines(), filename)
path: numba/tests/test_firstlinefinder.py
text: @user1:
```suggestion
        entry = (len(source), timestamp, source.splitlines(True), filename)
```
keepends for linecache

diff_hunk: @@ -79,6 +80,32 @@ def test_string_eval(self):
         # Cannot determine first line of string evaled functions
         self.assertIsNone(first_def_line)
 
+    def test_string_eval_linecache(self):
+        source = """def foo():
+            pass
+        """
path: numba/tests/test_firstlinefinder.py
text: @user2:
i am testing it with:
```python
        source = """if True:
        # indent causes IndentationError in ast.parse

        def decor(f):
            return f

        @user1 
        @user1
        @user1
        def foo():
            pass
        """
```
indentation is not handled. As `ast.parse()` is seeing:

```python
        @user1  
        @user1
        @user1
        def foo():
            pass
```

@author:
Thanks - I've added some tests now that do something similar.