Refactor kernel sim for the experimental kernel

numba_dpex/examples/kernel/kernel_simulator.py

@@ -0,0 +1,53 @@
+# SPDX-FileCopyrightText: 2023 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0
+
+import dpnp
+import numpy as np
+import numpy.testing as testing
+
+import numba_dpex as ndpx
+import numba_dpex.experimental as ndpx_ex
+
+
+def vector_sum(a, b, c):
+    i = ndpx.get_global_id(0)
+    c[i] = a[i] + b[i]
+
+
+# Utility function for printing and testing
+def driver(a, b, c, global_size):
+    # Sim result
+    c_sim = dpnp.zeros_like(c)
+
+    # Call sim kernel
+    ndpx_ex.call_kernel(vector_sum, ndpx.Range(global_size), a, b, c_sim)
+
+    # Call dpex kernel
+    ndpx_ex.call_kernel(
+        ndpx_ex.kernel(vector_sum), ndpx.Range(global_size), a, b, c
+    )
+
+    # Compare kernel result with simulator
+    testing.assert_equal(c.asnumpy(), c_sim.asnumpy())
+
+
+# Main function
+def main():
+    N = 10
+    global_size = N
+    print("Vector size N", N)
+
+    # Create random vectors on the default device
+    a = dpnp.random.random(N)
+    b = dpnp.random.random(N)
+    c = dpnp.ones_like(a)
+
+    print("Using device ...")
+    print(a.device)
+    driver(a, b, c, global_size)
+    print("Done...")
+
+
+if __name__ == "__main__":
+    main()

numba_dpex/experimental/launcher.py

@@ -33,6 +33,7 @@
     ItemType,
     NdItemType,
 )
+from numba_dpex.kernel_api.fake_kernel import FakeKernel
 
 
 class LLRange(NamedTuple):
@@ -218,8 +219,30 @@ def codegen(
     return sig, codegen
 
 
-@dpjit
 def call_kernel(kernel_fn, index_space, *kernel_args) -> None:
+    """Calls a numba_dpex.kernel decorated function.
+
+    Calls a python function or a `numba_dpex.experimental.kernel` decorated
+    function from CPython or from another dpjit function. Kernel execution
+    happens in synchronous way, therefore the thread will be blocked till the
+    kernel done execution.
+
+    Args:
+        kernel_fn (function | KernelDispatcher): A python function or a
+            `numba_dpex.kernel` decorated function that is compiled to a
+            `KernelDispatcher` by `numba_dpex`.
+        index_space (Range | NdRange): A `numba_dpex.Range` or
+            `numba_dpex.NdRange` type object that specifies the index space for
+            the kernel.
+    """
+    if isinstance(kernel_fn, SPIRVKernelDispatcher):
+        _call_kernel(kernel_fn, index_space, *kernel_args)
+    else:
+        FakeKernel(kernel_fn, index_space, *kernel_args).execute()
+
+
+@dpjit
+def _call_kernel(kernel_fn, index_space, *kernel_args) -> None:
     """Calls a numba_dpex.kernel decorated function from CPython or from another
     dpjit function. Kernel execution happens in synchronous way, so the thread
     will be blocked till the kernel done execution.
@@ -245,7 +268,7 @@ def call_kernel_async(
     kernel_fn,
     index_space,
     dependent_events: list[dpctl.SyclEvent],
-    *kernel_args
+    *kernel_args,
 ) -> tuple[dpctl.SyclEvent, dpctl.SyclEvent]:
     """Calls a numba_dpex.kernel decorated function from CPython or from another
     dpjit function. Kernel execution happens in asynchronous way, so the thread