Merge branch 'master' into refactor_interface

Author: isazi

CHANGELOG.md

@@ -4,6 +4,7 @@ This project adheres to [Semantic Versioning](http://semver.org/).
 
 ## Unreleased
 
+## [0.4.4] - 2023-03-09
 ### Added
 - Support for using time_limit in simulation mode
 - Helper functions for energy tuning
@@ -12,6 +13,7 @@ This project adheres to [Semantic Versioning](http://semver.org/).
 
 ### Changed
 - Changed what timings are stored in cache files
+- No longer inserting partial loop unrolling factor of 0 in CUDA
 
 ## [0.4.3] - 2022-10-19
 ### Added

doc/source/conf.py

@@ -59,9 +59,9 @@
 # built documents.
 #
 # The short X.Y version.
-version = u'0.4.3'
+version = u'0.4.4'
 # The full version, including alpha/beta/rc tags.
-release = u'0.4.3'
+release = u'0.4.4'
 
 # The language for content autogenerated by Sphinx. Refer to documentation
 # for a list of supported languages.

doc/source/optimization.rst

@@ -28,16 +28,26 @@ the ``strategy=`` optional argument of ``tune_kernel()``. Kernel Tuner currently
 
 Most strategies have some mechanism built in to detect when to stop tuning, which may be controlled through specific 
 parameters that can be passed to the strategies using the ``strategy_options=`` optional argument of ``tune_kernel()``. You 
-can also override whatever internal stop criterion the strategy uses, and set either a time limit in seconds or a maximum 
-number of unique function evaluations.
+can also override whatever internal stop criterion the strategy uses, and set either a time limit in seconds (using ``time_limit=``) or a maximum 
+number of unique function evaluations (using ``max_fevals=``).
+
+To give an example, one could simply add these two arguments to any code calling ``tune_kernel()``:
+
+.. code-block:: python
+
+    results, env = tune_kernel("vector_add", kernel_string, size, args, tune_params,
+                               strategy="random_sample",
+                               strategy_options=dict(max_fevals=5))
+
 
 A 'unique function evaluation' corresponds to the first time that Kernel Tuner tries to compile and benchmark a parameter 
 configuration that has been selected by the optimization strategy. If you are continuing from a previous tuning session using 
-cache files, serving a value from the cache also counts as a function evaluation for the strategy. Only unique function 
-evaluations are counted, so the second time a parameter configuration is selected by the strategy it is served from the 
+cache files, serving a value from the cache for the first time in the run also counts as a function evaluation for the strategy.
+Only unique function evaluations are counted, so the second time a parameter configuration is selected by the strategy it is served from the 
 cache, but not counted as a unique function evaluation.
 
-The ``strategy_options=`` argument of ``tune_kernel()`` should be used as follows:
+Below all the strategies are listed with their strategy-specific options that can be passed in a dictionary to the ``strategy_options=`` argument
+of ``tune_kernel()``.
 
 
 kernel_tuner.strategies.basinhopping

examples/c/matrix_multiply.cpp

@@ -0,0 +1,52 @@
+#!/usr/bin/env python
+"""
+    This is the vector_add example modified to show
+    how to use PythonKernel with the CuPy backend
+"""
+
+import cupy as cp
+import numpy as np
+from kernel_tuner.kernelbuilder import PythonKernel
+
+def kernelbuilder_example():
+
+    # To make this example self-contained we include the kernel as a string
+    # here, but you can also just point to a file with the kernel code
+    kernel_string = """
+    __global__ void vector_add(float *c, float *a, float *b, int n) {
+        int i = blockIdx.x * block_size_x + threadIdx.x;
+        if (i<n) {
+            c[i] = a[i] + b[i];
+        }
+    }
+    """
+
+    # Setup the arguments for our vector add kernel
+    size = 100000
+    a = cp.random.randn(size).astype(np.float32)
+    b = cp.random.randn(size).astype(np.float32)
+    c = cp.zeros_like(b)
+    n = np.int32(size)
+
+    # Note that the type and order should match our GPU code
+    # Because the arguments are all CuPy arrays, our PythonKernel does not need to
+    # worry about moving data between host and device
+    args = [c, a, b, n]
+
+    # We can instantiate a specific kernel configurations
+    params = {"block_size_x": 128}
+
+    # Here we construct a Python object that represents the kernel
+    # we can use it to conveniently use the GPU kernel in Python
+    # applications that want to frequently call the GPU kernel
+    vector_add = PythonKernel("vector_add", kernel_string, size, args, params, lang="cupy")
+
+    # We can use the PythonKernel instance as a regular Python function
+    vector_add(c, a, b, n)
+
+    # Compare the result in c with a+b computed in Python
+    assert np.allclose(c, a+b)
+
+
+if __name__ == "__main__":
+    kernelbuilder_example()

examples/c/matrix_multiply.py

@@ -54,7 +54,7 @@ __global__ void sum_floats(float *sum_global, floatvector *array, int n) {
         sum = sh_mem[ti];
         #pragma unroll
         for (unsigned int s=16; s>0; s>>=1) {
-            sum += __shfl_down_sync(0, sum, s);
+            sum += __shfl_down_sync(0xffffffff, sum, s);
         }
     }
     #else

examples/cuda/python_kernel_cupy.py

@@ -11,10 +11,10 @@ def tune():
 
     tune_params = OrderedDict()
     tune_params["block_size_x"] = [2**i for i in range(5,11)]
-    tune_params["use_shuffle"] = [0, 1]
     tune_params["vector"] = [2**i for i in range(3)]
     tune_params["num_blocks"] = [2**i for i in range(5,16)]
     tune_params["loop_unroll_factor_0"] = [0, 1, 8, 16, 32, 64]
+    tune_params["use_shuffle"] = [0, 1]
 
     problem_size = "num_blocks"
     size = 800000000

examples/cuda/reduction.cu

@@ -1,4 +1,4 @@
 from kernel_tuner.integration import store_results, create_device_targets
 from kernel_tuner.interface import tune_kernel, run_kernel
 
-__version__ = "0.4.3"
+__version__ = "0.4.4"