Merge pull request #1238 from IntelPython/fix/leftover_duplicate_code

Removes code duplication and adds new unit tests.

Author: Diptorup Deb

numba_dpex/core/codegen.py

@@ -2,7 +2,7 @@
 #
 # SPDX-License-Identifier: Apache-2.0
 
-import logging
+import warnings
 
 from llvmlite import binding as ll
 from llvmlite import ir as llvmir
@@ -31,7 +31,9 @@ def _optimize_functions(self, ll_module):
 
     @property
     def inline_threshold(self):
-        """The inlining threshold value to be used to optimize the final library"""
+        """
+        The inlining threshold value to be used to optimize the final library.
+        """
         if hasattr(self, "_inline_threshold"):
             return self._inline_threshold
         else:
@@ -41,11 +43,17 @@ def inline_threshold(self):
     def inline_threshold(self, value: int):
         """Returns the current inlining threshold level for the library."""
         if value < 0 or value > 3:
-            logging.warning(
+            warnings.warn(
                 "Unsupported inline threshold. Set a value between 0 and 3"
             )
             self._inline_threshold = 0
         else:
+            if value == 3:
+                warnings.warn(
+                    "Due to an existing compiler bug, setting INLINE_THRESHOLD "
+                    f"to {value} can lead to incorrect code generation on "
+                    "certain devices."
+                )
             self._inline_threshold = value
 
     def _optimize_final_module(self):
@@ -55,44 +63,20 @@ def _optimize_final_module(self):
         # Make optimization level depending on config.DPEX_OPT variable
         pmb.opt_level = config.DPEX_OPT
         if config.DPEX_OPT > 2:
-            logging.warning(
+            warnings.warn(
                 "Setting NUMBA_DPEX_OPT greater than 2 known to cause issues "
                 + "related to very aggressive optimizations that leads to "
                 + "broken code."
             )
 
         pmb.disable_unit_at_a_time = False
 
-        if config.INLINE_THRESHOLD is not None:
-            # Check if a decorator-level inline threshold was set and use that
-            # instead of the global configuration.
-            if (
-                hasattr(self, "_inline_threshold")
-                and self._inline_threshold > 0
-                and self._inline_threshold <= 3
-            ):
-                logging.warning(
-                    "Setting INLINE_THRESHOLD leads to very aggressive "
-                    + "optimizations that may produce incorrect binary."
-                )
-                pmb.inlining_threshold = self._inline_threshold
-            elif not hasattr(self, "_inline_threshold"):
-                logging.warning(
-                    "Setting INLINE_THRESHOLD leads to very aggressive "
-                    + "optimizations that may produce incorrect binary."
-                )
-                pmb.inlining_threshold = config.INLINE_THRESHOLD
-        else:
-            if (
-                hasattr(self, "_inline_threshold")
-                and self._inline_threshold > 0
-                and self._inline_threshold <= 3
-            ):
-                logging.warning(
-                    "Setting INLINE_THRESHOLD leads to very aggressive "
-                    + "optimizations that may produce incorrect binary."
-                )
-                pmb.inlining_threshold = self._inline_threshold
+        # The PassManagerBuilder's inlining_threshold property is set only when
+        # inline_threshold is g.t. 0. Doing otherwise, *i.e.*, setting the
+        # pmb.inlining_threshold to 0 will lead to at minimum `alwaysinline`
+        # pass to run.
+        if self.inline_threshold > 0:
+            pmb.inlining_threshold = self.inline_threshold
 
         pmb.disable_unroll_loops = True
         pmb.loop_vectorize = False

numba_dpex/core/kernel_interface/spirv_kernel.py

@@ -136,6 +136,15 @@ def compile(
         kernel = cres.target_context.prepare_spir_kernel(
             func, cres.signature.args
         )
+
+        # XXX: Setting the inline_threshold in the following way is a temporary
+        # workaround till the JitKernel dispatcher is replaced by
+        # experimental.dispatcher.KernelDispatcher.
+        if config.INLINE_THRESHOLD is not None:
+            cres.library.inline_threshold = config.INLINE_THRESHOLD
+        else:
+            cres.library.inline_threshold = 0
+
         cres.library._optimize_final_module()
         self._llvm_module = kernel.module.__str__()
         self._module_name = kernel.name

numba_dpex/tests/experimental/test_compiler_warnings.py

@@ -2,18 +2,33 @@
 #
 # SPDX-License-Identifier: Apache-2.0
 
-import warnings
-
+import dpctl
 import pytest
+from numba.core import types
 
 import numba_dpex as dpex
+from numba_dpex import DpctlSyclQueue, DpnpNdArray
 from numba_dpex import experimental as dpex_exp
+from numba_dpex import int64
+
+
+def _kernel(a, b, c):
+    i = dpex.get_global_id(0)
+    c[i] = a[i] + b[i]
 
 
 def test_compilation_mode_option_user_definition():
-    def kernel_func(a, b, c):
-        i = dpex.get_global_id(0)
-        c[i] = a[i] + b[i]
+    with pytest.warns(UserWarning):
+        dpex_exp.kernel(_compilation_mode="kernel")(_kernel)
+
+
+def test_inline_threshold_level_warning():
+    """
+    Test compiler warning generation with an inline_threshold value of 3.
+    """
 
-    with pytest.warns(warnings.warn(UserWarning)):
-        dpex_exp.kernel(_compilation_mode="kernel")(kernel_func)
+    with pytest.warns(UserWarning):
+        queue_ty = DpctlSyclQueue(dpctl.SyclQueue())
+        i64arr_ty = DpnpNdArray(ndim=1, dtype=int64, layout="C", queue=queue_ty)
+        kernel_sig = types.void(i64arr_ty, i64arr_ty, i64arr_ty)
+        dpex_exp.kernel(inline_threshold=3)(_kernel).compile(kernel_sig)

numba_dpex/tests/experimental/test_inline_threshold_config.py

@@ -34,7 +34,9 @@ def test_inline_threshold_set_using_config():
 
 
 def test_inline_threshold_set_using_decorator_option():
-    """Test setting the inline_threshold value using the kernel decorator flag"""
+    """
+    Test setting the inline_threshold value using the kernel decorator flag
+    """
 
     disp = dpex_exp.kernel(inline_threshold=2)(kernel_func)
     flags = compiler.Flags()

numba_dpex/tests/misc/test_warnings.py

@@ -1,6 +1,7 @@
-import logging
+import warnings
 
 import dpnp
+import pytest
 
 import numba_dpex as dpex
 import numba_dpex.config as config
@@ -11,32 +12,47 @@ def foo(a):
     a[dpex.get_global_id(0)] = 0
 
 
-def test_opt_warning(caplog):
+def test_opt_warning():
     bkp = config.DPEX_OPT
     config.DPEX_OPT = 3
 
-    with caplog.at_level(logging.WARNING):
+    with pytest.warns(UserWarning):
         foo[dpex.Range(10)](dpnp.arange(10))
 
     config.DPEX_OPT = bkp
 
-    assert "NUMBA_DPEX_OPT" in caplog.text
+
+def test_inline_threshold_eq_3_warning():
+    bkp = config.INLINE_THRESHOLD
+    config.INLINE_THRESHOLD = 3
+
+    with pytest.warns(UserWarning):
+        foo[dpex.Range(10)](dpnp.arange(10))
+
+    config.INLINE_THRESHOLD = bkp
 
 
-def test_inline_warning(caplog):
+def test_inline_threshold_negative_val_warning_():
     bkp = config.INLINE_THRESHOLD
-    config.INLINE_THRESHOLD = 2
+    config.INLINE_THRESHOLD = -1
 
-    with caplog.at_level(logging.WARNING):
+    with pytest.warns(UserWarning):
         foo[dpex.Range(10)](dpnp.arange(10))
 
     config.INLINE_THRESHOLD = bkp
 
-    assert "INLINE_THRESHOLD" in caplog.text
 
+def test_inline_threshold_gt_3_warning():
+    bkp = config.INLINE_THRESHOLD
+    config.INLINE_THRESHOLD = 4
 
-def test_no_warning(caplog):
-    with caplog.at_level(logging.WARNING):
+    with pytest.warns(UserWarning):
         foo[dpex.Range(10)](dpnp.arange(10))
 
-    assert caplog.text == ""
+    config.INLINE_THRESHOLD = bkp
+
+
+def test_no_warning():
+    with warnings.catch_warnings():
+        warnings.simplefilter("error")
+        foo[dpex.Range(10)](dpnp.arange(10))