Writing Device Functions

The user-level API of SYCL does not have a notion for device-only functions, i.e. functions that can be only invoked from a kernel and not from a host function. However, numba-dpex provides a special decorator numba_dpex.func specifically to implement device functions.

@ndpx.func
def a_device_function(a):
    return a + 1

To use a device function from an another device function:

@ndpx.func
def another_device_function(a):
    return a_device_function(a * 2)

To use a device function from a kernel function numba_dpex.kernel:

@ndpx.kernel
def a_kernel_function(a, b):
    i = ndpx.get_global_id(0)
    b[i] = another_device_function(a[i])

Unlike a kernel function, a device function can return a value like normal functions.

Todo

Specific capabilities and limitations for device functions need to be added.