Add DPEX_OPT, INLINE_THRESHOLD config and do not use numba's OPT

numba_dpex/config.py

@@ -87,3 +87,7 @@ def __getattr__(name):
 # Flag to turn on the ConstantSizeStaticLocalMemoryPass in the kernel pipeline.
 # The pass is turned off by default.
 STATIC_LOCAL_MEM_PASS = _readenv("NUMBA_DPEX_STATIC_LOCAL_MEM_PASS", int, 0)
+
+DPEX_OPT = _readenv("NUMBA_DPEX_OPT", int, 2)
+
+INLINE_THRESHOLD = _readenv("NUMBA_DPEX_INLINE_THRESHOLD", int, None)

numba_dpex/core/codegen.py

@@ -2,6 +2,8 @@
 #
 # SPDX-License-Identifier: Apache-2.0
 
+import logging
+
 from llvmlite import binding as ll
 from llvmlite import ir as llvmir
 from numba.core import utils
@@ -31,11 +33,22 @@ def _optimize_final_module(self):
         # Run some lightweight optimization to simplify the module.
         pmb = ll.PassManagerBuilder()
 
-        # Make optimization level depending on config.OPT variable
-        pmb.opt_level = config.OPT
+        # Make optimization level depending on config.DPEX_OPT variable
+        pmb.opt_level = config.DPEX_OPT
+        if config.DPEX_OPT > 2:
+            logging.warning(
+                "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
-        pmb.inlining_threshold = 2
+        if config.INLINE_THRESHOLD is not None:
+            logging.warning(
+                "Setting INLINE_THRESHOLD leads to very aggressive "
+                + "optimizations that may produce incorrect binary."
+            )
+            pmb.inlining_threshold = config.INLINE_THRESHOLD
         pmb.disable_unroll_loops = True
         pmb.loop_vectorize = False
         pmb.slp_vectorize = False

numba_dpex/core/kernel_interface/dispatcher.py

@@ -94,7 +94,7 @@ def __init__(
             self._kernel_bundle_cache = NullCache()
         self._cache_hits = 0
 
-        if debug_flags or config.OPT == 0:
+        if debug_flags or config.DPEX_OPT == 0:
             # if debug is ON we need to pass additional
             # flags to igc.
             self._create_sycl_kernel_bundle_flags = ["-g", "-cl-opt-disable"]

numba_dpex/core/parfors/kernel_builder.py

@@ -80,7 +80,7 @@ def _compile_kernel_parfor(
     )
 
     dpctl_create_program_from_spirv_flags = []
-    if debug or config.OPT == 0:
+    if debug or config.DPEX_OPT == 0:
         # if debug is ON we need to pass additional flags to igc.
         dpctl_create_program_from_spirv_flags = ["-g", "-cl-opt-disable"]
 

numba_dpex/tests/misc/test_warnings.py

@@ -0,0 +1,42 @@
+import logging
+
+import dpnp
+
+import numba_dpex as dpex
+import numba_dpex.config as config
+
+
+@dpex.kernel(enable_cache=False)
+def foo(a):
+    a[dpex.get_global_id(0)] = 0
+
+
+def test_opt_warning(caplog):
+    bkp = config.DPEX_OPT
+    config.DPEX_OPT = 3
+
+    with caplog.at_level(logging.WARNING):
+        foo[dpex.Range(10)](dpnp.arange(10))
+
+    config.DPEX_OPT = bkp
+
+    assert "NUMBA_DPEX_OPT" in caplog.text
+
+
+def test_inline_warning(caplog):
+    bkp = config.INLINE_THRESHOLD
+    config.INLINE_THRESHOLD = 2
+
+    with caplog.at_level(logging.WARNING):
+        foo[dpex.Range(10)](dpnp.arange(10))
+
+    config.INLINE_THRESHOLD = bkp
+
+    assert "INLINE_THRESHOLD" in caplog.text
+
+
+def test_no_warning(caplog):
+    with caplog.at_level(logging.WARNING):
+        foo[dpex.Range(10)](dpnp.arange(10))
+
+    assert caplog.text == ""