Huge overhead on devcloud linked to dpctl calls

[fcharras]

Version: numba_0.20.0dev3 and main

The three following dpctl calls 1 2 3 have huge wall time on edge devcloud (measured ranging from 10 to 30ms each call by py-spy, see speedscope report):

On the devcloud this add about 80 seconds to the k-means benchmark (for an expected 10 seconds).

I didn't see the issue on a local machine, but maybe the remaining small overhead that we reported comes from there.

@oleksandr-pavlyk not sure if this should be considered as an unreasonable use in numba_dpex (those calls should be expected to be that long and cached ?) or a bug in dpctl.

I've experimenting with caching the values and can confirm that caching those 3 calls completely remove the overhead.

Regarding the scope of the cache, I'll check if a hotfix that consists in storing those value in a WeakKeyDictionary where keys are val, and usm_mem, and wrapping SyclDevice(device) call in a lru_cache, is enough. (if so, will monkey-patch in sklearn_numba_dpex in the meantime).

[ogrisel]

To avoid confusion, the .svg file extension of the py-spy report file should be a .json extension. The .svg extension only make sense when py-spy is used to generate a flamegraph report as an SVG file instead of a json speedscope trace.

[ogrisel]

By zooming in the report, it seems that the overhead seems to come from the repeated calls to typeof_usm_ndarray:

but I cannot see the calls to dpctl.SyclDevice(device).

[fcharras]

The relevant calls to investigate here are the cells that are closer to the bottom, since it's as large as the parent cells, it means it's the bottleneck.

By hovering on those cells you can see the filename and the line number. You should be able to trace it back to the 3 lines I've linked in the OP.

[fcharras]

Workaround here soda-inria/sklearn-numba-dpex#104
edit: doesn't work
edit edit: this works https://github.com/soda-inria/sklearn-numba-dpex/blob/e040e78d2a5492d7b7b0ec79c2576f0df15cb9db/sklearn_numba_dpex/patches/load_numba_dpex.py#L44

[AlexanderKalistratov]

@ogrisel @fcharras could you please verify if #946 fixes the issue?

[fcharras]

Unfortunately, it doesn't fix. Looking at the PR it doesn't seem to change the instructions that lead to the time consuming steps in the OP (that are, dpctl.SyclDevice(device) and *.sycl_device.filter_string).

[fcharras]

The workaround I've posted yesterday doesn't work either. (currently fixing)

[oleksandr-pavlyk]

The dpctl.SyclDevice calls sycl::device constructor, which scores each available device and selects the one with the highest score. SyclDevice.filter_string calls sycl::get_devices() and searches for the given device in that list.

Construction of SYCL devices may thus be expensive, as RT must talk to the hardware. numba-dpex should not be constructing the device, but rather should capture it from the instance of usm_ndarray that it is inferring the type from. This is forthcoming, but I do not the ETA.

This suggests that using SYCL_DEVICE_FILTER to limit the number of devices discoverable by RT should improve the timing.
Use sycl-ls to determine the appropriate value to set the environment variable to. For example: with SYCL_DEVICE_FILTER=level_zero:gpu:0 the runtime would only discover one level-zero GPU device.

[AlexanderKalistratov]

@oleksandr-pavlyk

You are correct. We should extract device from usm_ndarray instead of creation of new one from filter_string.
But we still have to get filter_string since it is part of type signature. And getting filter_string is slow

[oleksandr-pavlyk]

I would argue that the need to store filter_string as part of type signature would be rendered unnecessary once boxing/unboxing of dpctl.SyclQueue is implemented.