Update User Guides about Debugging

README.md

@@ -66,7 +66,7 @@ python numba_dppy/examples/sum.py
 
 ## Debugging
 
-Please follow instructions in the [debugging.md](docs/user_guides/debugging.md)
+Please follow instructions in the [debugging](docs/user_guides/debugging.rst)
 
 ## Reporting issues
 

docs/conf.py

@@ -62,7 +62,7 @@
 # Add any paths that contain custom static files (such as style sheets) here,
 # relative to this directory. They are copied after the builtin static files,
 # so a file named "default.css" will overwrite the builtin "default.css".
-html_static_path = ["_static"]
+# html_static_path = ["_static"]
 
 modindex_common_prefix = ["numba_dppy."]
 

docs/index.rst

@@ -67,10 +67,10 @@ There are two ways to program SYCL devices using Numba-dppy:
    :maxdepth: 1
    :caption: User Guides
 
-    Getting Started <user_guides/getting_started.rst>
-    Programming SYCL Kernels <user_guides/kernel_programming_guide/index.rst>
-    Debugging with GDB <user_guides/debugging.md>
-    numba-dppy for numba.cuda Programmers <user_guides/migrating_from_numba_cuda.rst>
+    Getting Started <user_guides/getting_started>
+    Programming SYCL Kernels <user_guides/kernel_programming_guide/index>
+    Debugging with GDB <user_guides/debugging>
+    numba-dppy for numba.cuda Programmers <user_guides/migrating_from_numba_cuda>
 
 .. toctree::
     :maxdepth: 1

docs/user_guides/debugging.rst

@@ -0,0 +1,85 @@
+Debugging with GDB
+==================
+
+`numba-dppy` allows SYCL kernels to be debugged with the GDB debugger.
+Setting the debug environment variable :envvar:`NUMBA_DPPY_DEBUG` (e.g.
+:samp:`export NUMBA_DPPY_DEBUG=1`) enables the emission of debug information.
+To disable debugging, unset the variable, i.e :samp:`unset NUMBA_DPPY_DEBUG`.
+
+.. note::
+
+    Beware that enabling debug info significantly increases the memory consumption for each compiled kernel.
+    For large application, this may cause out-of-memory error.
+
+Not all GDB features supported by Numba on CPUs are yet supported in `numba-dppy`.
+Currently, the following debugging features are available:
+
+- Source location (filename and line number).
+- Setting break points by the line number.
+- Stepping over break points.
+
+.. note::
+
+    Debug features depend heavily on optimization level.
+    At full optimization (equivalent to O3), most of the variables are optimized out.
+    It is recommended to debug at "no optimization" level via :envvar:`NUMBA_OPT` (e.g. ``export NUMBA_OPT=0``).
+    For more information refer to the Numba documentation
+    `Debugging JIT compiled code with GDB <https://numba.pydata.org/numba-doc/latest/user/troubleshoot.html?highlight=numba_opt#debugging-jit-compiled-code-with-gdb>`_.
+
+Requirements
+------------
+
+Intel® Distribution for GDB is needed for `numba-dppy`'s debugging features
+to work. Intel® Distribution for GDB is part of Intel oneAPI and the relevant
+documentation can be found at
+`Intel® Distribution for GDB documentation <https://software.intel.com/content/www/us/en/develop/tools/oneapi/components/distribution-for-gdb.html>`_.
+
+Example of GDB usage
+--------------------
+
+For example, given the following `numba-dppy` kernel code:
+
+.. code-block:: python
+    :caption: sum.py
+    :linenos:
+
+    import numpy as np
+    import numba_dppy as dppy
+    import dpctl
+
+    @dppy.kernel
+    def data_parallel_sum(a, b, c):
+        i = dppy.get_global_id(0)
+        c[i] = a[i] + b[i]
+
+    global_size = 10
+    N = global_size
+
+    a = np.array(np.random.random(N), dtype=np.float32)
+    b = np.array(np.random.random(N), dtype=np.float32)
+    c = np.ones_like(a)
+
+    with dpctl.device_context("opencl:gpu") as gpu_queue:
+        data_parallel_sum[global_size, dppy.DEFAULT_LOCAL_SIZE](a, b, c)
+
+Running GDB and creating breakpoint in kernel:
+
+.. code-block:: bash
+
+    $ export NUMBA_DPPY_DEBUG=1
+    $ gdb-oneapi -q --args python sum.py
+    (gdb) break sum.py:7  # Set breakpoint in kernel
+    (gdb) run
+    Thread 2.2 hit Breakpoint 1,  with SIMD lanes [0-7], dppl_py_devfn__5F__5F_main_5F__5F__2E_data_5F_parallel_5F_sum_24_1_2E_array_28_float32_2C__20_1d_2C__20_C_29__2E_array_28_float32_2C__20_1d_2C__20_C_29__2E_array_28_float32_2C__20_1d_2C__20_C_29_ () at sum.py:7
+    7          i = dppy.get_global_id(0)
+    (gdb) next
+    8          c[i] = a[i] + b[i]
+
+Limitations
+-----------
+
+Currently, `numba-dppy` provides only initial support of debugging SYCL kernels.
+The following functionalities are **not supported**:
+
+- Printing kernel local variables (e.g. ``info locals``).
+- Stepping over several offloaded functions.

docs/user_guides/kernel_programming_guide/index.rst

@@ -1,14 +1,14 @@
-Programming SYCL Kernels Using ``numba_dppy.kernel``
-====================================================
+Programming SYCL Kernels Using :func:`~numba_dppy.decorators.kernel`
+====================================================================
 
 .. toctree::
    :maxdepth: 2
 
-   writing_kernels.rst
-   memory-management.rst
-   synchronization.rst
-   device-functions.rst
-   atomic-operations.rst
-   reduction.rst
-   ufunc.rst
-   supported-python-features.rst
+   writing_kernels
+   memory-management
+   synchronization
+   device-functions
+   atomic-operations
+   reduction
+   ufunc
+   supported-python-features