Modify test cases and examples to use level zero queue.

MANIFEST.in

@@ -2,6 +2,7 @@ include MANIFEST.in
 include README.md setup.py LICENSE
 
 recursive-include numba_dppy *.cl
+recursive-include numba_dppy *.spir
 
 include versioneer.py
 include numba_dppy/_version.py

numba_dppy/dpnp_glue/dpnp_array_creations_impl.py

@@ -64,11 +64,12 @@ def dpnp_zeros_like_impl(a, dtype=None):
     void dpnp_initval_c(void* result1, void* value, size_t size)
 
     """
-    res_dtype = dtype or a.dtype
-    if dtype:
-        name_dtype = res_dtype.dtype.name
-    else:
+    res_dtype = dtype
+    if dtype == types.none:
+        res_dtype = a.dtype
         name_dtype = res_dtype.name
+    else:
+        name_dtype = res_dtype.dtype.name
 
     sig = signature(ret_type, types.voidptr, types.voidptr, types.intp)
     dpnp_func = dpnp_ext.dpnp_func("dpnp_" + name, [name_dtype, "NONE"], sig)
@@ -77,7 +78,7 @@ def dpnp_zeros_like_impl(a, dtype=None):
 
     def dpnp_impl(a, dtype=None):
         b = np.zeros(1, dtype=res_dtype)
-        out = np.arange(a.size, dtype=res_dtype)
+        out = np.zeros(a.shape, dtype=res_dtype)
         common_impl(a, b, out, dpnp_func, PRINT_DEBUG)
         return out
 
@@ -98,11 +99,12 @@ def dpnp_ones_like_impl(a, dtype=None):
     void dpnp_initval_c(void* result1, void* value, size_t size)
 
     """
-    res_dtype = dtype or a.dtype
-    if dtype:
-        name_dtype = res_dtype.dtype.name
-    else:
+    res_dtype = dtype
+    if dtype == types.none:
+        res_dtype = a.dtype
         name_dtype = res_dtype.name
+    else:
+        name_dtype = res_dtype.dtype.name
 
     sig = signature(ret_type, types.voidptr, types.voidptr, types.intp)
     dpnp_func = dpnp_ext.dpnp_func("dpnp_" + name, [name_dtype, "NONE"], sig)
@@ -111,7 +113,7 @@ def dpnp_ones_like_impl(a, dtype=None):
 
     def dpnp_impl(a, dtype=None):
         b = np.ones(1, dtype=res_dtype)
-        out = np.arange(a.size, dtype=res_dtype)
+        out = np.ones(a.shape, dtype=res_dtype)
         common_impl(a, b, out, dpnp_func, PRINT_DEBUG)
         return out
 
@@ -139,13 +141,14 @@ def dpnp_full_like_impl(a, b):
     PRINT_DEBUG = dpnp_lowering.DEBUG
 
     def dpnp_impl(a, b):
-        out = np.arange(a.size, dtype=res_dtype)
+        out = np.ones(a.shape, dtype=res_dtype)
         common_impl(a, b, out, dpnp_func, PRINT_DEBUG)
         return out
 
     return dpnp_impl
 
 
+# TODO: This implementation is incorrect
 @overload(stubs.dpnp.full)
 def dpnp_full_impl(a, b):
     name = "full"

numba_dppy/dpnp_glue/dpnp_array_ops_impl.py

@@ -91,9 +91,13 @@ def dpnp_cumprod_impl(a):
     dpnp_func = dpnp_ext.dpnp_func("dpnp_" + name, [a.dtype.name, "NONE"], sig)
 
     PRINT_DEBUG = dpnp_lowering.DEBUG
+    if a.dtype == types.Integer:
+        ret_dtype = np.int64
+    else:
+        ret_dtype = a.dtype
 
     def dpnp_impl(a):
-        out = np.arange(a.size, dtype=a.dtype)
+        out = np.arange(a.size, dtype=ret_dtype)
         common_impl(a, out, dpnp_func, PRINT_DEBUG)
 
         return out

numba_dppy/dpnp_glue/dpnp_linalgimpl.py

@@ -346,11 +346,15 @@ def dpnp_matrix_power_impl(a, n):
                        size_t size_n, size_t size_k)
     """
 
+    PRINT_DEBUG = dpnp_lowering.DEBUG
+
     def dpnp_impl(a, n):
         if n < 0:
             raise ValueError("n < 0 is not supported for np.linalg.matrix_power(a, n)")
 
         if n == 0:
+            if PRINT_DEBUG:
+                print("dpnp implementation")
             return np.identity(a.shape[0], a.dtype)
 
         result = a

numba_dppy/dpnp_glue/dpnp_statisticsimpl.py

@@ -175,6 +175,7 @@ def dpnp_mean_impl(a):
         types.intp,
     )
     dpnp_func = dpnp_ext.dpnp_func("dpnp_" + name, [a.dtype.name, "NONE"], sig)
+    PRINT_DEBUG = dpnp_lowering.DEBUG
 
     res_dtype = np.float64
     if a.dtype == types.float32:

numba_dppy/dppy_passes.py

@@ -98,46 +98,58 @@ def run_pass(self, state):
                             ).value
                             if (
                                 isinstance(call_node, ir.Expr)
+                                and call_node.op == "getattr"
                                 and call_node.attr == "array"
                             ):
-                                arg = None
-                                # at first look in keyword arguments to get the shape, which has to be
-                                # constant
-                                if expr.kws:
-                                    for _arg in expr.kws:
-                                        if _arg[0] == "shape":
-                                            arg = _arg[1]
-
-                                if not arg:
-                                    arg = expr.args[0]
-
-                                error = False
-                                # arg can be one constant or a tuple of constant items
-                                arg_type = func_ir.get_definition(arg.name)
-                                if isinstance(arg_type, ir.Expr):
-                                    # we have a tuple
-                                    for item in arg_type.items:
+                                # let's check if it is from numba_dppy.local
+                                attr_node = block.find_variable_assignment(
+                                    call_node.value.name
+                                ).value
+                                if (
+                                    isinstance(attr_node, ir.Expr)
+                                    and attr_node.op == "getattr"
+                                    and attr_node.attr == "local"
+                                ):
+
+                                    arg = None
+                                    # at first look in keyword arguments to get the shape, which has to be
+                                    # constant
+                                    if expr.kws:
+                                        for _arg in expr.kws:
+                                            if _arg[0] == "shape":
+                                                arg = _arg[1]
+
+                                    if not arg:
+                                        arg = expr.args[0]
+
+                                    error = False
+                                    # arg can be one constant or a tuple of constant items
+                                    arg_type = func_ir.get_definition(arg.name)
+                                    if isinstance(arg_type, ir.Expr):
+                                        # we have a tuple
+                                        for item in arg_type.items:
+                                            if not isinstance(
+                                                func_ir.get_definition(item.name),
+                                                ir.Const,
+                                            ):
+                                                error = True
+                                                break
+
+                                    else:
                                         if not isinstance(
-                                            func_ir.get_definition(item.name), ir.Const
+                                            func_ir.get_definition(arg.name), ir.Const
                                         ):
                                             error = True
                                             break
 
-                                else:
-                                    if not isinstance(
-                                        func_ir.get_definition(arg.name), ir.Const
-                                    ):
-                                        error = True
-                                        break
-
-                                if error:
-                                    warnings.warn_explicit(
-                                        "The size of the Local memory has to be constant",
-                                        errors.NumbaError,
-                                        state.func_id.filename,
-                                        state.func_id.firstlineno,
-                                    )
-                                    raise
+                                    if error:
+                                        warnings.warn_explicit(
+                                            "The size of the Local memory has to be constant",
+                                            errors.NumbaError,
+                                            state.func_id.filename,
+                                            state.func_id.firstlineno,
+                                        )
+                                        raise
 
         if config.DEBUG or config.DUMP_IR:
             name = state.func_ir.func_id.func_qualname

numba_dppy/examples/blacksholes_kernel.py

@@ -0,0 +1,97 @@
+# Copyright 2021 Intel Corporation
+#
+# Licensed under the Apache License, Version 2.0 (the "License");
+# you may not use this file except in compliance with the License.
+# You may obtain a copy of the License at
+#
+#      http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing, software
+# distributed under the License is distributed on an "AS IS" BASIS,
+# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+# See the License for the specific language governing permissions and
+# limitations under the License.
+
+import numpy as np
+import math
+import time
+import numba_dppy, numba_dppy as dppy
+import unittest
+import dpctl
+
+
+RISKFREE = 0.02
+VOLATILITY = 0.30
+
+A1 = 0.31938153
+A2 = -0.356563782
+A3 = 1.781477937
+A4 = -1.821255978
+A5 = 1.330274429
+RSQRT2PI = 0.39894228040143267793994605993438
+
+
+def randfloat(rand_var, low, high):
+    return (1.0 - rand_var) * low + rand_var * high
+
+
+OPT_N = 400
+iterations = 2
+
+stockPrice = randfloat(np.random.random(OPT_N), 5.0, 30.0)
+optionStrike = randfloat(np.random.random(OPT_N), 1.0, 100.0)
+optionYears = randfloat(np.random.random(OPT_N), 0.25, 10.0)
+callResult = np.zeros(OPT_N)
+putResult = -np.ones(OPT_N)
+
+
+@dppy.kernel
+def black_scholes_dppy(callResult, putResult, S, X, T, R, V):
+    i = dppy.get_global_id(0)
+    if i >= S.shape[0]:
+        return
+    sqrtT = math.sqrt(T[i])
+    d1 = (math.log(S[i] / X[i]) + (R + 0.5 * V * V) * T[i]) / (V * sqrtT)
+    d2 = d1 - V * sqrtT
+
+    K = 1.0 / (1.0 + 0.2316419 * math.fabs(d1))
+    cndd1 = (
+        RSQRT2PI
+        * math.exp(-0.5 * d1 * d1)
+        * (K * (A1 + K * (A2 + K * (A3 + K * (A4 + K * A5)))))
+    )
+    if d1 > 0:
+        cndd1 = 1.0 - cndd1
+
+    K = 1.0 / (1.0 + 0.2316419 * math.fabs(d2))
+    cndd2 = (
+        RSQRT2PI
+        * math.exp(-0.5 * d2 * d2)
+        * (K * (A1 + K * (A2 + K * (A3 + K * (A4 + K * A5)))))
+    )
+    if d2 > 0:
+        cndd2 = 1.0 - cndd2
+
+    expRT = math.exp((-1.0 * R) * T[i])
+    callResult[i] = S[i] * cndd1 - X[i] * expRT * cndd2
+    putResult[i] = X[i] * expRT * (1.0 - cndd2) - S[i] * (1.0 - cndd1)
+
+
+blockdim = 512, 1
+griddim = int(math.ceil(float(OPT_N) / blockdim[0])), 1
+
+with dpctl.device_context("level0:gpu") as gpu_queue:
+    time1 = time.time()
+    for i in range(iterations):
+        black_scholes_dppy[blockdim, griddim](
+            callResult,
+            putResult,
+            stockPrice,
+            optionStrike,
+            optionYears,
+            RISKFREE,
+            VOLATILITY,
+        )
+
+print("callResult : ", callResult)
+print("putResult : ", putResult)

numba_dppy/examples/blacksholes_njit.py

@@ -2,7 +2,7 @@
 # Copyright (c) 2017 Intel Corporation
 # SPDX-License-Identifier: BSD-2-Clause
 #
-
+import dpctl
 import numba
 import numpy as np
 import math
@@ -16,7 +16,7 @@ def cndf2(inp):
     return out
 
 
-@numba.njit(parallel={"offload": True}, fastmath=True)
+@numba.njit(parallel=True, fastmath=True)
 def blackscholes(sptprice, strike, rate, volatility, timev):
     logterm = np.log(sptprice / strike)
     powterm = 0.5 * volatility * volatility
@@ -52,9 +52,24 @@ def main():
     args = parser.parse_args()
     options = args.options
 
-    run(10)
-    print("options = ", options)
-    run(options)
+    if dpctl.has_gpu_queues():
+        print("\nScheduling on OpenCL GPU\n")
+        with dpctl.device_context("opencl:gpu") as gpu_queue:
+            run(10)
+    else:
+        print("\nSkip scheduling on OpenCL GPU\n")
+    # if dpctl.has_gpu_queues(dpctl.backend_type.level_zero):
+    #    print("\nScheduling on Level Zero GPU\n")
+    #    with dpctl.device_context("level0:gpu") as gpu_queue:
+    #        run(10)
+    # else:
+    #    print("\nSkip scheduling on Level Zero GPU\n")
+    if dpctl.has_cpu_queues():
+        print("\nScheduling on OpenCL CPU\n")
+        with dpctl.device_context("opencl:cpu") as cpu_queue:
+            run(10)
+    else:
+        print("\nSkip scheduling on OpenCL CPU\n")
 
 
 if __name__ == "__main__":

numba_dppy/examples/pa_examples/test1-2d.py

@@ -32,7 +32,7 @@ def f1(a, b):
 print("a:", a, hex(a.ctypes.data))
 print("b:", b, hex(b.ctypes.data))
 
-with dpctl.device_context("opencl:gpu:0"):
+with dpctl.device_context("level0:gpu:0"):
     c = f1(a, b)
 
 print("BIG RESULT c:", c, hex(c.ctypes.data))