Regression: Unable to execute kernel on second Level-Zero device

[psalz]

Describe the bug

Executing a trivial kernel on the second out of two Level-Zero devices (Arc A770) in my machine causes a PI_ERROR_DEVICE_NOT_AVAILABLE error since #10794 was merged.

To Reproduce

The following program

#include <cstdio>
#include <sycl/sycl.hpp>

int main() {
    for(size_t i = 0; i < sycl::device::get_devices().size(); ++i) {
        auto device = sycl::device::get_devices()[i];
        printf("Using device %zu: %s (id: %u, platform: %s)\n", i,
            device.get_info<sycl::info::device::name>().c_str(),
            device.get_info<sycl::info::device::vendor_id>(),
            device.get_platform().get_info<sycl::info::platform::name>().c_str());
        sycl::queue q{device};
        q.parallel_for(sycl::range<1>(10), [](sycl::id<1>) {
            // no-op
        });
        q.wait_and_throw();
        printf("Done\n");
    }
}

runs for all devices except the last, which is the second Arc A770 in the system exposed through Level-Zero, for which it seems to hang for a couple of seconds and then crashes:

$ clang++ -fsycl test.cpp -o test && ./test
Using device 0: Intel(R) Xeon(R) Gold 5317 CPU @ 3.00GHz (id: 32902, platform: Intel(R) OpenCL)
Done
Using device 1: Intel(R) Arc(TM) A770 Graphics (id: 32902, platform: Intel(R) OpenCL Graphics)
Done
Using device 2: Intel(R) Arc(TM) A770 Graphics (id: 32902, platform: Intel(R) OpenCL Graphics)
Done
Using device 3: Intel(R) Arc(TM) A770 Graphics (id: 32902, platform: Intel(R) Level-Zero)
Done
Using device 4: Intel(R) Arc(TM) A770 Graphics (id: 32902, platform: Intel(R) Level-Zero)
terminate called after throwing an instance of 'sycl::_V1::runtime_error'
  what():  Native API failed. Native API returns: -2 (PI_ERROR_DEVICE_NOT_AVAILABLE) -2 (PI_ERROR_DEVICE_NOT_AVAILABLE)
[1]    439610 IOT instruction  ./test

It appears that this is not an issue with the second device per-se, but rather the runtime's handling of it. Point in case: If I limit the visible devices through ONEAPI_DEVICE_SELECTOR="level_zero:1" to only show the second device, everything still works:

$ ONEAPI_DEVICE_SELECTOR="level_zero:1" ./test
Using device 0: Intel(R) Arc(TM) A770 Graphics (id: 32902, platform: Intel(R) Level-Zero)
Done

I'm unfortunately unable to test whether this has been fixed since, as the current HEAD build of DPC++ (dbd9b67) segfaults during compilation. Edit: As noted below, current builds still have this problem.

Environment (please complete the following information):

• OS: Ubuntu 22.04
• Motherboard: Supermicro X12DPU-6
• Target device and vendor: 2x Intel Arc A770
• DPC++ version: clang version 17.0.0 (https://github.com/intel/llvm 0e49948)
• Dependencies version:

$ apt list --installed | grep "intel\|level"
intel-fw-gpu/jammy,jammy,now 2023.25.6-231~22.04 all [installed]
intel-gsc/jammy,now 0.8.9+51~u22.04 amd64 [installed,automatic]
intel-i915-dkms/jammy,jammy,now 1.23.5.19.230406.21.5.17.0.1034+i38-1 all [installed]
intel-igc-cm/jammy,now 1.0.176+i600~22.04 amd64 [installed]
intel-igc-core/now 1.0.14062.11 amd64 [installed,local]
intel-igc-opencl/now 1.0.14062.11 amd64 [installed,local]
intel-level-zero-gpu/now 1.3.26516.18 amd64 [installed,local]
intel-media-va-driver-non-free/jammy,now 23.2.1-647~22.04 amd64 [installed]
intel-metrics-discovery/jammy,now 1.12.164-647~22.04 amd64 [installed,automatic]
intel-metrics-library/jammy,now 1.0.133-647~22.04 amd64 [installed,automatic]
intel-microcode/jammy-updates,jammy-security,now 3.20230808.0ubuntu0.22.04.1 amd64 [installed,automatic]
intel-opencl-icd/now 23.22.26516.18 amd64 [installed,local]
intel-platform-cse-dkms/jammy,now 2023.11.1-36 amd64 [installed]
intel-platform-vsec-dkms/jammy,now 2023.20.0-21 amd64 [installed]
level-zero-devel/now 1.11.0 amd64 [installed,local]
level-zero/now 1.11.0 amd64 [installed,upgradable to: 1.11.0-647~22.04]
libdrm-intel1/jammy-updates,now 2.4.113-2~ubuntu0.22.04.1 amd64 [installed,automatic]

[AlexeySachkov]

@jandres742: FYI

[abagusetty]

I was running the same example on a device with multiple PVCs using this compiler (9249bf8) and runtime 23.22.26516.18.

The behavior seems to be hung for the second device and last couple of calls to trace is shown below.

Using device 0: Intel(R) Data Center GPU Max 1550 (id: 32902, platform: Intel(R) Level-Zero)
Done
Using device 1: Intel(R) Data Center GPU Max 1550 (id: 32902, platform: Intel(R) Level-Zero)

The trace shown below:

---> piEventRelease(
	pi_event : 0x2ed2510
UR ---> urEventReleaseInternal(Event)
UR <--- urEventReleaseInternal(Event)(UR_RESULT_SUCCESS)
) ---> 	pi_result : PI_SUCCESS

---> piQueueFinish(
	<unknown> : 0x2d929e0
UR ---> EventCreate(Queue->Context, Queue, HostVisible.value(), Event)
UR <--- EventCreate(Queue->Context, Queue, HostVisible.value(), Event)(UR_RESULT_SUCCESS)
UR ---> urEventRetain(*Event)
UR <--- urEventRetain(*Event)(UR_RESULT_SUCCESS)

[al42and]

Can confirm, the problem happens with a8ab2bb and 0e49948, but not with the commit before.

$ clang++ -fsycl 10982.cpp -o test && ONEAPI_DEVICE_SELECTOR=level_zero:gpu ./test
Using device 0: Intel(R) Arc(TM) A770 Graphics (id: 32902, platform: Intel(R) Level-Zero)
Done
Using device 1: Intel(R) UHD Graphics 770 (id: 32902, platform: Intel(R) Level-Zero)
terminate called after throwing an instance of 'sycl::_V1::runtime_error'
  what():  Native API failed. Native API returns: -999 (Unknown PI error) -999 (Unknown PI error)
Aborted (core dumped)

$ ONEAPI_DEVICE_SELECTOR=level_zero:gpu SYCL_PI_TRACE=-1 ./test
[....]
---> piEnqueueKernelLaunch(
        <unknown> : 0x3254d60
        <unknown> : 0x28b5aa0
        <unknown> : 1
        <unknown> : 0x7ffe4fc6fb58
        <unknown> : 0x7ffe4fc6fb28
        <unknown> : 0
        <unknown> : 0
        pi_event * : 0[ nullptr ]
        pi_event * : 0x28b1748[ 0 ... ]
UR ---> TmpWaitList.createAndRetainUrZeEventList( NumEventsInWaitList, EventWaitList, Queue, UseCopyEngine)
UR <--- TmpWaitList.createAndRetainUrZeEventList( NumEventsInWaitList, EventWaitList, Queue, UseCopyEngine)(UR_RESULT_SUCCESS)
UR ---> Queue->Context->getAvailableCommandList( Queue, CommandList, UseCopyEngine, true )
UR ---> insertStartBarrierIfDiscardEventsMode(CommandList)
UR <--- insertStartBarrierIfDiscardEventsMode(CommandList)(UR_RESULT_SUCCESS)
UR ---> insertActiveBarriers(CommandList, UseCopyEngine)
UR <--- insertActiveBarriers(CommandList, UseCopyEngine)(UR_RESULT_SUCCESS)
UR <--- Queue->Context->getAvailableCommandList( Queue, CommandList, UseCopyEngine, true )(UR_RESULT_SUCCESS)
UR ---> createEventAndAssociateQueue(Queue, Event, UR_COMMAND_KERNEL_LAUNCH, CommandList, IsInternal)
UR ---> EventCreate(Queue->Context, Queue, HostVisible.value(), Event)
UR <--- EventCreate(Queue->Context, Queue, HostVisible.value(), Event)(UR_RESULT_SUCCESS)
UR ---> urEventRetain(*Event)
UR <--- urEventRetain(*Event)(UR_RESULT_SUCCESS)
UR <--- createEventAndAssociateQueue(Queue, Event, UR_COMMAND_KERNEL_LAUNCH, CommandList, IsInternal)(UR_RESULT_SUCCESS)
UR ---> urKernelRetain(Kernel)
UR <--- urKernelRetain(Kernel)(UR_RESULT_SUCCESS)
UR ---> Queue->executeCommandList(CommandList, false, true)
UR <--- Queue->executeCommandList(CommandList, false, true)(UR_RESULT_SUCCESS)
) --->  pi_result : PI_SUCCESS
        [out]pi_event * : 0[ nullptr ]
        [out]pi_event * : 0x28b1748[ 0x35ac320 ... ]

---> piEventRelease(
        pi_event : 0x35ac320
UR ---> urEventReleaseInternal(Event)
UR <--- urEventReleaseInternal(Event)(UR_RESULT_SUCCESS)
) --->  pi_result : PI_SUCCESS

---> piQueueFinish(
        <unknown> : 0x3254d60
UR ---> UrQueue->executeAllOpenCommandLists()
UR <--- UrQueue->executeAllOpenCommandLists()(UR_RESULT_ERROR_UNKNOWN)
) --->  pi_result : -999

terminate called after throwing an instance of 'sycl::_V1::runtime_error'
  what():  Native API failed. Native API returns: -999 (Unknown PI error) -999 (Unknown PI error)
Aborted (core dumped)

Using the per-device contexts avoids the problem:

#include <cstdio>
#include <sycl/sycl.hpp>

int main() {
    for(size_t i = 0; i < sycl::device::get_devices().size(); ++i) {
        auto device = sycl::device::get_devices()[i];
        sycl::context ctx{device};
        printf("Using device %zu: %s (id: %u, platform: %s)\n", i,
            device.get_info<sycl::info::device::name>().c_str(),
            device.get_info<sycl::info::device::vendor_id>(),
            device.get_platform().get_info<sycl::info::platform::name>().c_str());
        sycl::queue q{ctx, device};
        q.parallel_for(sycl::range<1>(10), [](sycl::id<1>) {
            // no-op
        });
        q.wait_and_throw();
        printf("Done\n");
    }
}

[github-actions]

This issue is stale because it has been open 180 days with no activity. Remove stale label or comment or this will be automatically closed in 30 days.

[psalz]

It appears this has been fixed by the switch to the unified runtime, @abagusetty, @al42and can you confirm?

[abagusetty]

Can be closed, fixed on PVC using [1.3.30049]

[al42and]

Can confirm, works fine with 9eb93e0 / 1.3.30049.