intel / intel/llvm

[L0][CUDA][HIP] Event dependency bug

Open
#8,132 12 comments 0 reactions 0 assignees View on GitHub
bug cuda hip
Dominant language
LLVM
Stars
1.5k
Forks
854
Avg merge
3d 17h
Merged PRs (30d)
137

Description

**Description**
I have ran into an issue with level zero's event behaviour.
I have found that an event which depends on another event is completed before it should.
It is unclear if this is an issue with the event info or the kernel is not waiting on the prior event.
There is a reproducer below. I have also included some hardware information. Let me know if you require any additional details.
The problem only exists with the level zero backend and does not occur with OpenCL.

**Reproducer**
The code below creates two tasks. The first copies memory from the host to the device, the second executes a basic kernel.
The code uses USM. We create a dependency between the fist task and the second.

The events status are checked to make sure the second task waits on the first before executing.
To do this the depender status is checked in a while loop until it is either running or complete.
Then once the first task is running or finished the second task is checked to see if it has completed yet.

If it does an error is created. We would expect that each backend would not throw the error.

```
#include

void check_dependency(cl::sycl::event dependee_e, cl::sycl::event depender_e) {
cl::sycl::info::event_command_status depender_status =
depender_e.get_info();

while (depender_status != cl::sycl::info::event_command_status::running &&
depender_status != cl::sycl::info::event_command_status::complete) {
depender_status =
depender_e.get_info();
}
if (dependee_e.get_info() !=
cl::sycl::info::event_command_status::complete) {
throw std::runtime_error("Oh no!");
}
}

int main(int argc, char const *argv[]) {
sycl::queue q{};
constexpr int n = 10'000'000;
float *d_A = sycl::malloc_device(n, q);
float *h_A = sycl::malloc_host(n, q);

constexpr int n_repeat = 1000;

for (int i = 0; i < n_repeat; ++i) {
auto ev1 = q.memcpy(d_A, h_A, n * sizeof(float));
auto ev2 = q.submit([&](sycl::handler &cgh) {
cgh.depends_on({ev1});
cgh.parallel_for(sycl::nd_range<1>(32, 32), [=](sycl::nd_item<1> item) {
int x = item.get_global_linear_id();
d_A[x] = x;
});
});
check_dependency(ev1, ev2);
q.wait_and_throw();
}

std::cout << "Success\n";

sycl::free(d_A, q);
sycl::free(h_A, q);
return 0;
}
```

Results:
Level zero GPU
```
$ SYCL_DEVICE_FILTER=level_zero:gpu ./a.out
terminate called after throwing an instance of 'std::runtime_error'
what(): Oh no!
Aborted (core dumped)
```

OpenCL GPU
```
$ SYCL_DEVICE_FILTER=opencl:gpu ./a.out
Success
```

OpenCL CPU
```
$ SYCL_DEVICE_FILTER=opencl:cpu ./a.out
Success
```

**Environment:**

- OS:
```
Distributor ID: Ubuntu
Description: Ubuntu 20.04.5 LTS
Release: 20.04
Codename: focal
```
- DPC++ version:
```
clang version 16.0.0 (https://github.com/AidanBeltonS/llvm/ 1fda2a37568b67ad929c703e0078dbe645e002cb)
Target: x86_64-unknown-linux-gnu
Thread model: posix
InstalledDir: /home/aidan/llvm/build/bin
```
- sycl-ls output:
[opencl:cpu:0] Intel(R) OpenCL, Intel(R) Core(TM) i7-10610U CPU @ 1.80GHz 3.0 [2022.13.3.0.16_160000]
[opencl:cpu:1] Intel(R) OpenCL, Intel(R) Core(TM) i7-10610U CPU @ 1.80GHz 3.0 [2022.13.3.0.16_160000]
[opencl:gpu:2] Intel(R) OpenCL HD Graphics, Intel(R) UHD Graphics [0x9b41] 3.0 [21.36.20889]
[ext_oneapi_level_zero:gpu:0] Intel(R) Level-Zero, Intel(R) UHD Graphics [0x9b41] 1.1 [1.2.20889]

Contributor guide

Open the contributing guide

Assessment

This issue has not been assessed yet.

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.