Gaps in Writing Device Functions

[PokhodenkoSA]

Study the Writing Device Functions section and document the equivalent feature in numba-dppy.

Features:

1. Device function

@cuda.jit(device=True)
def a_device_function(a, b):
...

numba-dppy provides dppy.func:

@dppy.func
def bar(a):
    return a*a

@dppy.kernel
def foo(in, out):
    i = dppy.get_global_id(0)
    out[i] = bar(in[i])

Similar:

• invoked from within the device (by a kernel or another device function)
• return a value (kernels can not return values)

Questions about differences:

• dppy.func cannot be invoked from the host. cuda.jit(device=True) too?
• Decorator supports type specialization as a kernel. cuda?
• Implementation details?
• dppy.func is inlined into the OpenCL Program containing the @dppy.kernel function calling it. SYCL implementation the same as OpenCL ?
• Examples from numba.cuda for it?

Examples:

• https://github.com/IntelPython/numba-dppy/blob/main/numba_dppy/examples/dppy_func.py

Missing features:

• Rename dppy.kernel and dppy.func to dppy.jit #189 dppy.jit
• Support for Calling Device Functions from numba.jit functions #205 Call from @numba.jit
• Support for Calling Device Functions from Ufuncs #193 Call from @vectorize