update documentation

doc/source/backends.rst

@@ -1,6 +1,7 @@
 .. toctree::
    :maxdepth: 2
 
+.. _backends:
 
 Backends
 ========
@@ -15,24 +16,30 @@ language, but sometimes you'll want to specifically choose a backend.
 CUDA Backends
 -------------
 
-PyCUDA is default CUDA backend in Kernel Tuner. It is comparable in feature completeness with CuPy.
+Kernel Tuner automatically selects a CUDA backend when it detects CUDA as the language of the
+kernel to be tuned. Depending on whether the required dependencies are installed it first checks
+if it could use the ``cuda-python`` backend, then ``cupy``, then PyCUDA. This behavior can be overwritten
+by manually selecting a backend with the ``lang=`` option. 
 Because the HIP kernel language is identical to the CUDA kernel language, HIP is included here as well.
-To use HIP on nvidia GPUs, see https://github.com/jatinx/hip-on-nv.
+To use HIP on nvidia GPUs, see https://github.com/jatinx/hip-on-nv. HIP must be indicated manually
+using the ``lang=`` option.
 
-While the PyCUDA backend expects all inputs and outputs to be Numpy arrays, the CuPy backend also 
+Kernel inputs and outputs can be handeld slightly differently by different backends, but the default
+way to pass arguments to the kernel is through Numpy arrays, which is supported by all backends.
+For example, while the PyCUDA backend expects all inputs and outputs to be Numpy arrays, the CuPy backend also 
 supports cupy arrays as input and output arguments for the kernels. This gives the user more control 
 over how memory is handled by Kernel Tuner. Also checks during output verification can happen 
 entirely on the GPU when using only cupy arrays.
 
-Texture memory is only supported by the PyCUDA backend, while the CuPy backend is the only one that 
-support C++ signatures for the kernels. With the other backends, it is required that the kernel has 
+Texture memory is only supported by the PyCUDA backend. However, some other limitations apply to the PyCUDA
+backend, for example support for kernels with C++ signatures. PyCUDA requires that the kernel has 
 extern "C" linkage. If not, the entire code is wrapped in an extern "C" block, which may cause issues 
 if the code also contains C++ code that cannot have extern "C" linkage, including code that may be 
 present in header files.
 
-As detailed further in :ref:`templates`, templated kernels are fully supported by the CuPy backend and 
-limited support is implemented by Kernel Tuner to support templated kernels for the PyCUDA and 
-CUDA-Python backends.
+Templated kernels and kernels with C++ signatures are supported by the cuda-python, CuPy, and HIP backends. 
+As detailed further in :ref:`templates`, Kernel Tuner has limited support for templated kernels
+when using the PyCUDA backend.
 
 
 .. csv-table:: Backend feature support
@@ -45,13 +52,17 @@ CUDA-Python backends.
   Constant memory,        ✓,  ✓,  ✓,  ✓
   Dynamic shared memory,  ✓,  ✓,  ✓,  ✓
   Texture memory,         ✓,  ✗,  ✗,  ✗
-  C++ kernel signature,   ✗,  ✓,  ✗,  ✗
-  Templated kernels,      ✓,  ✓,  ✓,  ✗
+  C++ kernel signature,   ✗,  ✓,  ✓,  ✓
+  Templated kernels,      ✓,  ✓,  ✓,  ✓
 
 
 Another important difference between the different backends is the compiler that is used. The table 
 below lists which Python package is required, how the backend can be selected and which compiler is 
-used to compile the kernels.
+used to compile the kernels. Note that backends relying on ``nvrtc`` can be more strict about
+including headers and also having host code in the code that is passed to the compiler. When using
+``cuda-python`` it is generally recommended to use an environment variable called ``CUDA_HOME`` that
+points to the root directory of the CUDA installation. For example, by adding 
+``export CUDA_HOME=/usr/local/cuda`` to your ``.bashrc`` file.
 
 
 .. csv-table:: Backend usage and compiler

doc/source/templates.rst

@@ -6,9 +6,17 @@
 Templated kernels
 -----------------
 
-It is quite common in CUDA programming to write kernels that use C++ templates. This can be very useful when writing code that can work for several types, for example floats and doubles. However, the use of C++ templates makes it slightly more difficult to directly 
-integrate the CUDA kernel into applications that are not written in C++, for example Matlab, Fortran, or Python. And since Kernel Tuner is written in Python, we needed to take a few extra steps to provide support for templated CUDA kernels. Let's first look at an 
-example of what it's like to tune a templated kernel with Kernel Tuner.
+It is quite common in CUDA programming to write kernels that use C++ templates. This can be very useful when writing code that can 
+work for several types, for example floats and doubles. However, the use of C++ templates makes it slightly more difficult to 
+directly integrate the CUDA kernel into applications that are not written in C++, for example Matlab, Fortran, or Python. And 
+since Kernel Tuner is written in Python, we needed to take a few extra steps to provide support for templated CUDA kernels.
+
+Templated kernels and kernels with C++ signatures are now fully supported when using the cuda-python, cupy, and HIP backends. 
+Kernel Tuner implements some limited support for templated kernels when using the PyCUDA backend. If you use templated kernels, it is
+recommended that you use either the cuda-python or cupy backend.
+
+The rest of this section explains how to use Kernel Tuner with templated kernels.
+Let's first look at an example of what it's like to tune a templated kernel with Kernel Tuner.
 
 Example
 ~~~~~~~
@@ -46,35 +54,32 @@ Then the Python script to tune this kernel would be as follows:
 
     tune_kernel("vector_add<float>", "vector_add.cu", size, args, tune_params)
 
-What you can see is that in the Python code we specify the template instantiation to use. Kernel Tuner will detect the use of templated kernels when the kernel_name positional argument to tune_kernel contains a template argument.
+What you can see is that in the Python code we specify the template instantiation to use.
+Kernel Tuner will detect the use of templated kernels when the ``kernel_name`` positional
+argument to ``tune_kernel`` contains a template argument.
 
-This feature also allows use to auto-tune template parameters to the kernel. We could for example define a tunable parameter:
+This feature also allows users to auto-tune template parameters of the templated kernel.
+We could for example define a tunable parameter:
 
 .. code-block:: python
 
     tune_params["my_type"] = ["float", "double"]
 
-and call tune_kernel using a tunable parameter inside the template arguments:
+and call ``tune_kernel`` using a tunable parameter inside the template arguments:
 
 .. code-block:: python
 
     tune_kernel("vector_add<my_type>", "vector_add.cu", size, args, tune_params)
 
-Selecting a backend
-~~~~~~~~~~~~~~~~~~~
-
-Kernel Tuner supports multiple backends, for CUDA these are based on PyCUDA and Cupy. The following explains how to enable tuning of templated kernels with either backend.
-
-The PyCuda backend is the default backend in Kernel Tuner and is selected if the user does not supply the 'lang' option and CUDA code is detected in the kernel source, or when lang is set to "CUDA" by the user. PyCuda requires CUDA kernels to have extern C linkage, 
-which means that C++ templated kernels are not supported. To support templated kernels regardless of this limitation Kernel Tuner attempts to wrap the templated CUDA kernel by inserting a compile-time template instantiation statement and a wrapper kernel that calls 
-the templated CUDA kernel, which is actually demoted to a __device__ function in the process. These automatic code rewrites have a real risk of breaking the code. To minimize the chance of errors due to Kernel Tuner's automatic code rewrites, it's best to isolate the 
-templated kernel in a single source file and include it where needed in the larger application.
-
-The Cupy backend provides much more advanced support for C++ templated kernels, because it internally uses NVRTC, the Nvidia runtime compiler. NVRTC does come with some restrictions however, for example NVRTC does not allow any host code to be inside code that
-is passed. So, like with the PyCuda backend it helps to separate the source code of device and host functions into seperate files. You can force Kernel Tuner to use the Cupy backend by passing the lang="cupy" option to tune_kernel. 
-
-
-
-
-
 
+Support for templated kernels is supported in most backends. However, support for templates is 
+limited when using the PyCUDA backend. PyCuda requires CUDA kernels to have extern C linkage, 
+which means that C++ templated kernels are not supported by PyCUDA directly. To support 
+templated kernels regardless of this limitation Kernel Tuner attempts to wrap the templated 
+CUDA kernel by inserting a compile-time template instantiation statement and a wrapper kernel 
+that calls the templated CUDA kernel, which is actually demoted to a __device__ function in 
+the process. These automatic code rewrites have a real risk of breaking the code. It is 
+therefore recommended to use the cuda-python or cupy backends instead. However, if you must 
+use PyCUDA, to minimize the chance of errors due to Kernel Tuner's automatic code rewrites, 
+it's best to isolate the templated kernel in a single source file and include it where needed 
+in the larger application.

kernel_tuner/searchspace.py

@@ -67,6 +67,7 @@ def __init__(
             adjacent: picks closest parameter value in both directions for each parameter
             Hamming: any parameter config with 1 different parameter value is a neighbor
             Hamming-adjacent: differs by closest parameter value for exactly 1 parameter.
+
         Optionally sort the searchspace by the order in which the parameter values were specified. By default, sort goes from first to last parameter, to reverse this use sort_last_param_first.
         Optionally an imported cache can be used instead with `from_cache`, in which case the `tune_params`, `restrictions` and `max_threads` arguments can be set to None, and construction is skipped.
         Optionally construction can be deffered to a later time by setting `defer_construction` to True, in which case the searchspace is not built on instantiation (experimental).