Merge branch 'main' into c_vecextensions_target

.github/workflows/ci.yml

@@ -95,16 +95,18 @@ jobs:
         -   uses: actions/checkout@v2
         -   name: "Main Script"
             run: |
-                EXTRA_INSTALL="matplotlib ipykernel nbconvert"
                 curl -L -O https://tiker.net/ci-support-v0
                 . ./ci-support-v0
                 build_py_project_in_conda_env
 
+                # ipython_genutils for https://github.com/jupyter/nbconvert/issues/1725
+                with_echo pip install matplotlib ipykernel nbconvert ipython_genutils
+
                 install_ispc
 
                 export PYOPENCL_TEST=portable:pthread
 
-                . ./build-py-project-and-run-examples.sh
+                . ./.ci/examples-funcs.sh
                 run_py_examples
                 run_ipynb_examples
                 run_floopy_examples

.gitlab-ci.yml

@@ -112,16 +112,19 @@ Pytest POCL Examples:
   stage: test
   script: |
     export PYOPENCL_TEST=portable:pthread
-    export EXTRA_INSTALL="pybind11 numpy mako pyvisfile matplotlib ipykernel nbconvert"
+    export EXTRA_INSTALL="pybind11 numpy mako"
 
     curl -L -O https://tiker.net/ci-support-v0
     . ./ci-support-v0
 
     build_py_project_in_venv
 
+    # ipython_genutils for https://github.com/jupyter/nbconvert/issues/1725
+    pip install matplotlib ipykernel nbconvert ipython_genutils
+
     install_ispc
 
-    . ./build-py-project-and-run-examples.sh
+    . ./.ci/examples-funcs.sh
     run_py_examples
     run_ipynb_examples
     run_floopy_examples

.pylintrc-local.yml

@@ -1,6 +1,16 @@
 - arg: extension-pkg-whitelist
   val: islpy
+
 - arg: ignore
-  val: compyte
+  val:
+  - compyte
+
+- arg: ignored-modules
+  val:
+  - IPython
+  - pycuda
+  - maptlotlib
+  - maptlotlib.pyplot
+
 - arg: init-hook
   val: sys.setrecursionlimit(5000)

.test-conda-env-py3.yml

@@ -11,8 +11,3 @@ dependencies:
 - mako
 - pyopencl
 - islpy
-
-# Only needed to make pylint succeed
-- matplotlib-base
-- ipykernel
-- ply

doc/tutorial.rst

@@ -560,7 +560,7 @@ Consider this example:
     #define lid(N) ((int) get_local_id(N))
     ...
       for (int i_outer = 0; i_outer <= -1 + (15 + n) / 16; ++i_outer)
-        for (int i_inner = 0; i_inner <= (-16 + n + -16 * i_outer >= 0 ? 15 : -1 + n + -16 * i_outer); ++i_inner)
+        for (int i_inner = 0; i_inner <= (-17 + n + -16 * i_outer >= 0 ? 15 : -1 + n + -16 * i_outer); ++i_inner)
           a[16 * i_outer + i_inner] = 0.0f;
     ...
 

loopy/check.py

@@ -1453,7 +1453,7 @@ def _are_sub_array_refs_equivalent(sar1, sar2, caller):
 
     from loopy.symbolic import SubstitutionMapper
     from pymbolic.mapper.substitutor import make_subst_func
-    from loopy.isl_helpers import simplify_via_aff
+    from loopy.symbolic import simplify_via_aff
     subst_func = make_subst_func({iname1.name:  iname2
                                   for iname1, iname2 in zip(sar1.swept_inames,
                                                             sar2.swept_inames)

loopy/isl_helpers.py

@@ -1,6 +1,5 @@
 """isl helpers"""
 
-
 __copyright__ = "Copyright (C) 2012 Andreas Kloeckner"
 
 __license__ = """
@@ -28,6 +27,7 @@
 
 import islpy as isl
 from islpy import dim_type
+from warnings import warn
 
 
 def pw_aff_to_aff(pw_aff):
@@ -448,19 +448,6 @@ def add_in_dims(aff):
 # }}}
 
 
-def simplify_via_aff(expr):
-    from loopy.symbolic import aff_to_expr, guarded_aff_from_expr, get_dependencies
-    from loopy.diagnostic import ExpressionToAffineConversionError
-
-    deps = sorted(get_dependencies(expr))
-    try:
-        return aff_to_expr(guarded_aff_from_expr(
-            isl.Space.create_from_names(isl.DEFAULT_CONTEXT, list(deps)),
-            expr))
-    except ExpressionToAffineConversionError:
-        return expr
-
-
 def project_out(set, inames):
     for iname in inames:
         var_dict = set.get_var_dict()
@@ -849,4 +836,11 @@ def find_and_rename_dim(isl_obj, dt, old_name, new_name):
 # }}}
 
 
+def simplify_via_aff(expr):
+    warn("simplify_via_aff has moved to loopy.symbolic. "
+            "Importing it from loopy.isl_helpers will stop working in July 2022.",
+            DeprecationWarning, stacklevel=2)
+    from loopy.symbolic import simplify_via_aff
+    return simplify_via_aff(expr)
+
 # vim: foldmethod=marker

loopy/library/reduction.py

@@ -188,8 +188,15 @@ def get_le_neutral(dtype):
         elif dtype.numpy_dtype.itemsize == 8:
             # 64 bit integer
             return var("LONG_MAX")
-    else:
-        raise NotImplementedError("less")
+    elif dtype.numpy_dtype.kind == "u":
+        if dtype.numpy_dtype.itemsize == 4:
+            # 32 bit integer
+            return var("UINT_MAX")
+        elif dtype.numpy_dtype.itemsize == 8:
+            # 64 bit integer
+            return var("ULONG_MAX")
+
+    raise NotImplementedError(f"neutral element for <= and {dtype}")
 
 
 def get_ge_neutral(dtype):
@@ -211,8 +218,10 @@ def get_ge_neutral(dtype):
         elif dtype.numpy_dtype.itemsize == 8:
             # 64 bit integer
             return var("LONG_MIN")
-    else:
-        raise NotImplementedError("less")
+    elif dtype.numpy_dtype.kind == "u":
+        return 0
+
+    raise NotImplementedError(f"neutral element for >= and {dtype}")
 
 
 class MaxReductionOperation(ScalarReductionOperation):

loopy/match.py

@@ -228,6 +228,10 @@ def __init__(self, tag: pytools.tag.Tag):
     def __call__(self, kernel, matchable):
         return self.tag in matchable.tags
 
+    def update_persistent_hash(self, key_hash, key_builder):
+        key_builder.rec(key_hash, type(self).__name__)
+        key_builder.rec(key_hash, self.tag)
+
 
 class GlobMatchExpressionBase(MatchExpressionBase):
     def __init__(self, glob):

loopy/symbolic.py

@@ -1834,6 +1834,11 @@ def map_reduction(self, expr):
         raise TypeError("reduction in '%s' not supported "
                 "for as-pwaff evaluation" % expr)
 
+    def map_call(self, expr):
+        # FIXME: There are some things here that we could handle, e.g. "abs".
+        raise TypeError(f"call in '{expr}' not supported "
+                "for as-pwaff evaluation")
+
 
 def aff_from_expr(space, expr, vars_to_zero=None):
     if vars_to_zero is None:
@@ -1957,14 +1962,29 @@ def qpolynomial_from_expr(space, expr):
 
 # {{{ simplify using aff
 
+def simplify_via_aff(expr):
+    from loopy.symbolic import aff_to_expr, guarded_aff_from_expr, get_dependencies
+    from loopy.diagnostic import ExpressionToAffineConversionError
+
+    deps = sorted(get_dependencies(expr))
+    try:
+        return aff_to_expr(guarded_aff_from_expr(
+            isl.Space.create_from_names(isl.DEFAULT_CONTEXT, list(deps)),
+            expr))
+    except ExpressionToAffineConversionError:
+        return expr
+
+
 @memoize_on_first_arg
 def simplify_using_aff(kernel, expr):
     """
     Simplifies *expr* on *kernel*'s domain.
 
     :arg expr: An instance of :class:`pymbolic.primitives.Expression`.
     """
-    inames = get_dependencies(expr) & kernel.all_inames()
+    deps = get_dependencies(expr)
+
+    inames = deps & kernel.all_inames()
 
     # FIXME: Ideally, we should find out what inames are usable and allow
     # the simplification to use all of those. For now, fall back to making
@@ -1974,6 +1994,15 @@ def simplify_using_aff(kernel, expr):
             .get_inames_domain(inames)
             .project_out_except(inames, [dim_type.set]))
 
+    non_inames = deps - set(domain.get_var_dict().keys())
+    non_inames = {name for name in set(non_inames) if name.isidentifier()}
+    if non_inames:
+        cur_dim = domain.dim(isl.dim_type.set)
+        domain = domain.insert_dims(isl.dim_type.set, cur_dim, len(non_inames))
+        for non_iname in sorted(non_inames):
+            domain = domain.set_dim_name(isl.dim_type.set, cur_dim, non_iname)
+            cur_dim += 1
+
     try:
         aff = guarded_aff_from_expr(domain.space, expr)
     except ExpressionToAffineConversionError:

loopy/target/c/codegen/expression.py

@@ -332,9 +332,13 @@ def map_integer_div_operator(self, base_func_name, op_func, expr, type_context):
         assumptions, domain = isl.align_two(assumption_non_param, domain)
         domain = domain & assumptions
 
+        num_type = self.infer_type(expr.numerator)
+        den_type = self.infer_type(expr.denominator)
         from loopy.isl_helpers import is_nonnegative
-        num_nonneg = is_nonnegative(expr.numerator, domain)
-        den_nonneg = is_nonnegative(expr.denominator, domain)
+        num_nonneg = is_nonnegative(expr.numerator, domain) \
+            or num_type.numpy_dtype.kind == "u"
+        den_nonneg = is_nonnegative(expr.denominator, domain) \
+            or den_type.numpy_dtype.kind == "u"
 
         result_dtype = self.infer_type(expr)
         suffix = result_dtype.numpy_dtype.type.__name__

loopy/target/opencl.py

@@ -189,7 +189,33 @@ class OpenCLCallable(ScalarCallable):
     def with_types(self, arg_id_to_dtype, callables_table):
         name = self.name
 
-        # unary functions
+        # {{{ unary functions
+        if name == "abs":
+            for id in arg_id_to_dtype:
+                if not -1 <= id <= 0:
+                    raise LoopyError(f"'{name}' can take only one argument.")
+
+            if 0 not in arg_id_to_dtype or arg_id_to_dtype[0] is None:
+                return (
+                        self.copy(arg_id_to_dtype=arg_id_to_dtype),
+                        callables_table)
+
+            dtype = arg_id_to_dtype[0].numpy_dtype
+
+            if dtype.kind in ("u", "i"):
+                # OpenCL C 2.2, Section 6.13.3: abs returns *u*gentype
+                from loopy.types import to_unsigned_dtype
+                return (self.copy(name_in_target=name,
+                            arg_id_to_dtype={
+                                0: NumpyType(dtype),
+                                -1: NumpyType(to_unsigned_dtype(dtype))}),
+                        callables_table)
+            elif dtype.kind == "f":
+                name = "fabs"
+            else:
+                raise LoopyTypeError(f"'{name}' does not support type {dtype}")
+
+        # deliberately not elif: abs branch above may end up taking this.
         if name in ["fabs", "acos", "asin", "atan", "cos", "cosh", "sin", "sinh",
                     "tan", "tanh", "exp", "log", "log10", "sqrt", "ceil", "floor",
                     "erf", "erfc"]:
@@ -199,8 +225,6 @@ def with_types(self, arg_id_to_dtype, callables_table):
                     raise LoopyError(f"'{name}' can take only one argument.")
 
             if 0 not in arg_id_to_dtype or arg_id_to_dtype[0] is None:
-                # the types provided aren't mature enough to specialize the
-                # callable
                 return (
                         self.copy(arg_id_to_dtype=arg_id_to_dtype),
                         callables_table)
@@ -219,6 +243,9 @@ def with_types(self, arg_id_to_dtype, callables_table):
                         arg_id_to_dtype={0: NumpyType(dtype), -1:
                             NumpyType(dtype)}),
                     callables_table)
+
+        # }}}
+
         # binary functions
         elif name in ["fmax", "fmin", "atan2", "copysign"]:
 
@@ -231,8 +258,6 @@ def with_types(self, arg_id_to_dtype, callables_table):
 
             if 0 not in arg_id_to_dtype or 1 not in arg_id_to_dtype or (
                     arg_id_to_dtype[0] is None or arg_id_to_dtype[1] is None):
-                # the types provided aren't mature enough to specialize the
-                # callable
                 return (
                         self.copy(arg_id_to_dtype=arg_id_to_dtype),
                         callables_table)

loopy/target/pyopencl.py

@@ -89,7 +89,7 @@ def with_types(self, arg_id_to_dtype, callables_table):
         if name in ["sqrt", "exp", "log",
                 "sin", "cos", "tan",
                 "sinh", "cosh", "tanh",
-                "conj", "abs"]:
+                "conj"]:
             if dtype.is_complex():
                 # function parameters are complex.
                 if dtype.numpy_dtype == np.complex64:
@@ -103,21 +103,15 @@ def with_types(self, arg_id_to_dtype, callables_table):
                         self.copy(name_in_target=f"{tpname}_{name}",
                             arg_id_to_dtype={0: dtype, -1: dtype}),
                         callables_table)
-            else:
-                # function calls for floating-point parameters.
-                numpy_dtype = dtype.numpy_dtype
-                if numpy_dtype.kind in ("u", "i"):
-                    dtype = NumpyType(np.float32)
-                if name == "abs":
-                    name = "fabs"
-                return (
-                        self.copy(name_in_target=name,
-                            arg_id_to_dtype={0: dtype, -1: dtype}),
-                        callables_table)
 
-        return (
-                self.copy(arg_id_to_dtype=arg_id_to_dtype),
-                callables_table)
+            # fall back to pure OpenCL for real-valued arguments
+
+        from loopy.target.opencl import OpenCLCallable
+        return OpenCLCallable(name,
+                arg_id_to_dtype=self.arg_id_to_dtype,
+                arg_id_to_descr=self.arg_id_to_descr,
+                name_in_target=self.name_in_target).with_types(
+                        arg_id_to_dtype, callables_table)
 
     def generate_preambles(self, target):
         name = self.name_in_target

loopy/tools.py

@@ -922,7 +922,8 @@ def _get_persistent_hashable_arg(arg):
             if args and isinstance(args[0], LoopKernel):
                 proc_log_str = f"{func.__name__} on '{args[0].name}'"
             elif args and isinstance(args[0], TranslationUnit):
-                proc_log_str = f"{func.__name__} on '{args[0].entrypoints}'"
+                entrypoints_str = ", ".join(args[0].entrypoints)
+                proc_log_str = f"{func.__name__} on '{entrypoints_str}'"
             else:
                 proc_log_str = f"{func.__name__}"
 

loopy/transform/buffer.py

@@ -102,7 +102,7 @@ def map_array_access(self, index, expn_state):
                 continue
 
             ax_index = index[i]
-            from loopy.isl_helpers import simplify_via_aff
+            from loopy.symbolic import simplify_via_aff
             ax_index = simplify_via_aff(
                     ax_index - abm.storage_base_indices[i])