With context

numba/core/cpu_dispatcher.py

@@ -6,22 +6,8 @@ class CPUDispatcher(dispatcher.Dispatcher):
     targetdescr = cpu_target
 
     def __init__(self, py_func, locals={}, targetoptions={}, impl_kind='direct', pipeline_class=compiler.Compiler):
-        if ('parallel' in targetoptions and isinstance(targetoptions['parallel'], dict) and
-                'offload' in targetoptions['parallel'] and  targetoptions['parallel']['offload'] == True):
-            import numba.dppl_config as dppl_config
-            if dppl_config.dppl_present:
-                from numba.dppl.compiler import DPPLCompiler
-                dispatcher.Dispatcher.__init__(self, py_func, locals=locals,
-                        targetoptions=targetoptions, impl_kind=impl_kind, pipeline_class=DPPLCompiler)
-            else:
-                print("---------------------------------------------------------------------------")
-                print("WARNING : offload=True option ignored. Ensure OpenCL drivers are installed.")
-                print("---------------------------------------------------------------------------")
-                dispatcher.Dispatcher.__init__(self, py_func, locals=locals,
-                    targetoptions=targetoptions, impl_kind=impl_kind, pipeline_class=pipeline_class)
-        else:
-            dispatcher.Dispatcher.__init__(self, py_func, locals=locals,
-                targetoptions=targetoptions, impl_kind=impl_kind, pipeline_class=pipeline_class)
+        dispatcher.Dispatcher.__init__(self, py_func, locals=locals,
+            targetoptions=targetoptions, impl_kind=impl_kind, pipeline_class=pipeline_class)
 
 
 dispatcher_registry['cpu'] = CPUDispatcher

numba/core/cpu_options.py

@@ -45,7 +45,6 @@ class ParallelOptions(object):
     Options for controlling auto parallelization.
     """
     def __init__(self, value):
-        self.gen_spirv = False
         if isinstance(value, bool):
             self.enabled = value
             self.comprehension = value
@@ -64,7 +63,6 @@ def __init__(self, value):
             self.stencil = value.pop('stencil', True)
             self.fusion = value.pop('fusion', True)
             self.prange = value.pop('prange', True)
-            self.gen_spirv = value.pop('offload', False)
             if value:
                 msg = "Unrecognized parallel options: %s" % value.keys()
                 raise NameError(msg)

numba/core/decorators.py

@@ -11,6 +11,8 @@
 from numba.core.errors import DeprecationError, NumbaDeprecationWarning
 from numba.stencils.stencil import stencil
 from numba.core import config, sigutils, registry, cpu_dispatcher
+from numba.dppl import gpu_dispatcher
+import dppl.ocldrv as ocldrv
 
 
 _logger = logging.getLogger(__name__)
@@ -175,16 +177,28 @@ def bar(x, y):
         return wrapper
 
 
+def get_current_disp(target):
+    gpu_env = None
+    if ocldrv.runtime.has_gpu_device():
+        gpu_env = ocldrv.runtime.get_gpu_device().get_env_ptr()
+    current_env = None
+    if ocldrv.runtime.has_current_device():
+        current_env = ocldrv.runtime.get_current_device().get_env_ptr()
+    if (gpu_env == current_env) and ocldrv.is_available():
+        return registry.dispatcher_registry['dppl']
+
+    return registry.dispatcher_registry[target]
+
+
 def _jit(sigs, locals, target, cache, targetoptions, **dispatcher_args):
-    dispatcher = registry.dispatcher_registry[target]
 
-    def wrapper(func):
+    def wrapper(func, dispatcher):
         if config.ENABLE_CUDASIM and target == 'cuda':
             from numba import cuda
             return cuda.jit(func)
         if config.DISABLE_JIT and not target == 'npyufunc':
             return func
-        if target == 'dppl':
+        if target == 'dppl_kernel':
             from . import dppl
             return dppl.jit(func)
         disp = dispatcher(py_func=func, locals=locals,
@@ -202,7 +216,13 @@ def wrapper(func):
                 disp.disable_compile()
         return disp
 
-    return wrapper
+    def __wrapper(func):
+        func.__compiled = {}
+        disp = get_current_disp(target)
+        if not disp in func.__compiled.keys():
+            func.__compiled[disp] = wrapper(func, disp)
+        return func.__compiled[disp]
+    return __wrapper
 
 
 def generated_jit(function=None, target='cpu', cache=False,

numba/dppl/__init__.py

@@ -226,7 +226,7 @@ def main():
     from numba import njit
     import numpy as np
 
-    @njit(parallel={'offload':True})
+    @njit(target='dppl')
     def f1(a, b):
         c = a + b
         return c

numba/dppl/examples/blacksholes_njit.py

@@ -14,7 +14,7 @@ def cndf2(inp):
     out = 0.5 + 0.5 * math.erf((math.sqrt(2.0)/2.0) * inp)
     return out
 
-@numba.njit(parallel={'offload':True}, fastmath=True)
+@numba.njit(target='dppl', fastmath=True)
 def blackscholes(sptprice, strike, rate, volatility, timev):
     logterm = np.log(sptprice / strike)
     powterm = 0.5 * volatility * volatility

numba/dppl/examples/pa_examples/test1-2d.py

@@ -1,7 +1,7 @@
 from numba import njit, gdb
 import numpy as np
 
-@njit(parallel={'offload':True})
+@njit(target='dppl')
 def f1(a, b):
     c = a + b
     return c

numba/dppl/examples/pa_examples/test1-3d.py

@@ -1,7 +1,7 @@
 from numba import njit, gdb
 import numpy as np
 
-@njit(parallel={'offload':True})
+@njit(target='dppl')
 def f1(a, b):
     c = a + b
     return c

numba/dppl/examples/pa_examples/test1-4d.py

@@ -1,7 +1,7 @@
 from numba import njit, gdb
 import numpy as np
 
-@njit(parallel={'offload':True})
+@njit(target='dppl')
 def f1(a, b):
     c = a + b
     return c

numba/dppl/examples/pa_examples/test1-5d.py

@@ -1,7 +1,7 @@
 from numba import njit, gdb
 import numpy as np
 
-@njit(parallel={'offload':True})
+@njit(target='dppl')
 def f1(a, b):
     c = a + b
     return c

numba/dppl/examples/pa_examples/test1.py

@@ -2,7 +2,7 @@
 import numpy as np
 
 
-@njit(parallel={'offload':True})
+@njit(target='dppl')
 def f1(a, b):
     c = a + b
     return c

numba/dppl/gpu_dispatcher.py

@@ -0,0 +1,23 @@
+from numba.core import dispatcher, compiler
+from numba.core.registry import cpu_target, dispatcher_registry
+import numba.dppl_config as dppl_config
+from numba.dppl.compiler import DPPLCompiler
+from numba.core.cpu_options import ParallelOptions
+
+
+class GPUDispatcher(dispatcher.Dispatcher):
+    targetdescr = cpu_target
+
+    def __init__(self, py_func, locals={}, targetoptions={}, impl_kind='direct', pipeline_class=compiler.Compiler):
+        if dppl_config.dppl_present:
+            targetoptions['parallel'] = True
+            dispatcher.Dispatcher.__init__(self, py_func, locals=locals,
+                    targetoptions=targetoptions, impl_kind=impl_kind, pipeline_class=DPPLCompiler)
+        else:
+            print("---------------------------------------------------------------------")
+            print("WARNING : DPPL pipeline ignored. Ensure OpenCL drivers are installed.")
+            print("---------------------------------------------------------------------")
+            dispatcher.Dispatcher.__init__(self, py_func, locals=locals,
+                targetoptions=targetoptions, impl_kind=impl_kind, pipeline_class=pipeline_class)
+
+dispatcher_registry['dppl'] = GPUDispatcher

numba/dppl/initialize.py

@@ -10,7 +10,7 @@ def init_jit():
 
 def initialize_all():
     from numba.core.registry import dispatcher_registry
-    dispatcher_registry.ondemand['dppl'] = init_jit
+    dispatcher_registry.ondemand['dppl_kernel'] = init_jit
 
     ll.load_library_permanently(find_library('DPPLOpenCLInterface'))
     ll.load_library_permanently(find_library('OpenCL'))

numba/dppl/tests/dppl/test_dppl_fallback.py

@@ -40,7 +40,7 @@ def np_rand_fallback():
             return a
 
         def run_dppl():
-            dppl = numba.njit(parallel={'offload':True})(np_rand_fallback)
+            dppl = numba.njit(target='dppl')(np_rand_fallback)
             return dppl()
 
         ref = np_rand_fallback

numba/dppl/tests/dppl/test_numpy_bit_twiddling_functions.py

@@ -13,7 +13,7 @@
 
 class TestNumpy_bit_twiddling_functions(DPPLTestCase):
     def test_bitwise_and(self):
-        @njit(parallel={'offload':True})
+        @njit(target='dppl')
         def f(a, b):
             c = np.bitwise_and(a, b)
             return c
@@ -27,7 +27,7 @@ def f(a, b):
 
 
     def test_bitwise_or(self):
-        @njit(parallel={'offload':True})
+        @njit(target='dppl')
         def f(a, b):
             c = np.bitwise_or(a, b)
             return c
@@ -41,7 +41,7 @@ def f(a, b):
 
 
     def test_bitwise_xor(self):
-        @njit(parallel={'offload':True})
+        @njit(target='dppl')
         def f(a, b):
             c = np.bitwise_xor(a, b)
             return c
@@ -55,7 +55,7 @@ def f(a, b):
 
 
     def test_bitwise_not(self):
-        @njit(parallel={'offload':True})
+        @njit(target='dppl')
         def f(a):
             c = np.bitwise_not(a)
             return c
@@ -68,7 +68,7 @@ def f(a):
 
 
     def test_invert(self):
-        @njit(parallel={'offload':True})
+        @njit(target='dppl')
         def f(a):
             c = np.invert(a)
             return c
@@ -81,7 +81,7 @@ def f(a):
 
 
     def test_left_shift(self):
-        @njit(parallel={'offload':True})
+        @njit(target='dppl')
         def f(a, b):
             c = np.left_shift(a, b)
             return c
@@ -95,7 +95,7 @@ def f(a, b):
 
 
     def test_right_shift(self):
-        @njit(parallel={'offload':True})
+        @njit(target='dppl')
         def f(a, b):
             c = np.right_shift(a, b)
             return c

numba/dppl/tests/dppl/test_numpy_comparison_functions.py

@@ -14,7 +14,7 @@ class TestNumpy_comparison_functions(DPPLTestCase):
     a = np.array([4,5,6])
     b = np.array([2,6,6])
     def test_greater(self):
-        @njit(parallel={'offload':True})
+        @njit(target='dppl')
         def f(a, b):
             c = np.greater(a, b)
             return c
@@ -25,7 +25,7 @@ def f(a, b):
 
 
     def test_greater_equal(self):
-        @njit(parallel={'offload':True})
+        @njit(target='dppl')
         def f(a, b):
             c = np.greater_equal(a, b)
             return c
@@ -36,7 +36,7 @@ def f(a, b):
 
 
     def test_less(self):
-        @njit(parallel={'offload':True})
+        @njit(target='dppl')
         def f(a, b):
             c = np.less(a, b)
             return c
@@ -47,7 +47,7 @@ def f(a, b):
 
 
     def test_less_equal(self):
-        @njit(parallel={'offload':True})
+        @njit(target='dppl')
         def f(a, b):
             c = np.less_equal(a, b)
             return c
@@ -58,7 +58,7 @@ def f(a, b):
 
 
     def test_not_equal(self):
-        @njit(parallel={'offload':True})
+        @njit(target='dppl')
         def f(a, b):
             c = np.not_equal(a, b)
             return c
@@ -69,7 +69,7 @@ def f(a, b):
 
 
     def test_equal(self):
-        @njit(parallel={'offload':True})
+        @njit(target='dppl')
         def f(a, b):
             c = np.equal(a, b)
             return c
@@ -80,7 +80,7 @@ def f(a, b):
 
 
     def test_logical_and(self):
-        @njit(parallel={'offload':True})
+        @njit(target='dppl')
         def f(a, b):
             c = np.logical_and(a, b)
             return c
@@ -94,7 +94,7 @@ def f(a, b):
 
 
     def test_logical_or(self):
-        @njit(parallel={'offload':True})
+        @njit(target='dppl')
         def f(a, b):
             c = np.logical_or(a, b)
             return c
@@ -108,7 +108,7 @@ def f(a, b):
 
 
     def test_logical_xor(self):
-        @njit(parallel={'offload':True})
+        @njit(target='dppl')
         def f(a, b):
             c = np.logical_xor(a, b)
             return c
@@ -122,7 +122,7 @@ def f(a, b):
 
 
     def test_logical_not(self):
-        @njit(parallel={'offload':True})
+        @njit(target='dppl')
         def f(a):
             c = np.logical_not(a)
             return c
@@ -135,7 +135,7 @@ def f(a):
 
 
     def test_maximum(self):
-        @njit(parallel={'offload':True})
+        @njit(target='dppl')
         def f(a, b):
             c = np.maximum(a, b)
             return c
@@ -149,7 +149,7 @@ def f(a, b):
 
 
     def test_minimum(self):
-        @njit(parallel={'offload':True})
+        @njit(target='dppl')
         def f(a, b):
             c = np.minimum(a, b)
             return c
@@ -163,7 +163,7 @@ def f(a, b):
 
 
     def test_fmax(self):
-        @njit(parallel={'offload':True})
+        @njit(target='dppl')
         def f(a, b):
             c = np.fmax(a, b)
             return c
@@ -177,7 +177,7 @@ def f(a, b):
 
 
     def test_fmin(self):
-        @njit(parallel={'offload':True})
+        @njit(target='dppl')
         def f(a, b):
             c = np.fmin(a, b)
             return c