get_global_id does not work if function is directly exported

[geexie]

If I try the following code

import dpctl
import base_l2_distance
import math

from numba_dppy import kernel, get_global_id, atomic, DEFAULT_LOCAL_SIZE
atomic_add = atomic.add

@kernel(access_types={"read_only": ["a", "b"], "write_only": ["c"]})
def l2_distance_kernel(a, b, c):
    i = get_global_id(0)
    j = get_global_id(1)
    sub = a[i, j] - b[i, j]
    sq = sub**2
    atomic_add(c, 0, sq)

def l2_distance(a, b, distance):
    with dpctl.device_context(base_l2_distance.get_device_selector()):
        l2_distance_kernel[(a.shape[0], a.shape[1]), DEFAULT_LOCAL_SIZE](a, b, distance)

    return math.sqrt(distance)

the following error is risen during typeinfer

  File "/nfs/site/home/mkolpako/code/miniconda3/envs/gpu_dev/lib/python3.7/site-packages/numba/core/typeinfer.py", line 1466, in resolve_value_type
    raise TypingError(msg, loc=inst.loc)
numba.core.errors.TypingError: Failed in dppy_nopython mode pipeline (step: nopython frontend)
Untyped global name 'get_global_id': Cannot determine Numba type of <class 'function'>

File "l2_distance_kernel.py", line 59:
def l2_distance_kernel(a, b, c):
    i = get_global_id(0)

Everything works if I use numba_dppy.get_global_id notation

[PokhodenkoSA]

Environment:
dpctl 0.11.1
numba 0.54.1
numba-dppy 0.15.0

All examples shows using this functions only via module.

[PokhodenkoSA]

Addressed in #694.

[diptorupd]

Letting users write get_global_id instead of numba_dpex.get_global_id is a syntactic preference IMO. If we want to be more descriptive perhaps we can go the route of numba.cuda and introduce a sycl module and change all of these functions to sycl module.

from numba_dpex import sycl

@sycl.kernel
def vecadd(a,b,c):
i = sycl.get_global_id(0)
a[i] = b[i] + c[i]

Having the module qualifier is much better in terms of overall readability and letting users know that the code is specific to the numba_dpex kernel API and will not work outside of the kernel function or without the kernel decorator. I am opening a separate issue #795 to track that work.

Supporting these intrinsic functions directly without using the module name is not going to be added at this moment.