[OFFLOAD] Improve handling of synchronization errors in L0 plugin and reenable tests#186927
Open
[OFFLOAD] Improve handling of synchronization errors in L0 plugin and reenable tests#186927
Conversation
|
✅ With the latest revision this PR passed the C/C++ code formatter. |
Contributor
Author
Member
|
@llvm/pr-subscribers-offload Author: None (fineg74) ChangesThis change improves handling of errors during synchronization in Level Zero plugin by ensuring cleanup of queues and events in case of an synchronization error. As a result multiple tests stopped hanging. Patch is 23.08 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/186927.diff 40 Files Affected:
diff --git a/offload/plugins-nextgen/level_zero/include/L0Device.h b/offload/plugins-nextgen/level_zero/include/L0Device.h
index 001a41ba77d7b..ab80cd505f1fe 100644
--- a/offload/plugins-nextgen/level_zero/include/L0Device.h
+++ b/offload/plugins-nextgen/level_zero/include/L0Device.h
@@ -651,5 +651,19 @@ class L0DeviceTy final : public GenericDeviceTy {
interop_spec_t *Prefers) override;
};
+struct L0DeviceQueueGuard {
+ L0DeviceQueueGuard(AsyncQueueTy *Queue, L0DeviceTy *Device)
+ : queue(Queue), device(Device) {}
+ ~L0DeviceQueueGuard() {
+ for (auto &Event : queue->WaitEvents)
+ auto Err = device->releaseEvent(Event);
+ queue->WaitEvents.clear();
+ }
+
+private:
+ AsyncQueueTy *queue;
+ L0DeviceTy *device;
+};
+
} // namespace llvm::omp::target::plugin
#endif // OPENMP_LIBOMPTARGET_PLUGINS_NEXTGEN_LEVEL_ZERO_L0DEVICE_H
diff --git a/offload/plugins-nextgen/level_zero/src/L0Device.cpp b/offload/plugins-nextgen/level_zero/src/L0Device.cpp
index ef521ab681b4f..821f4d9222138 100644
--- a/offload/plugins-nextgen/level_zero/src/L0Device.cpp
+++ b/offload/plugins-nextgen/level_zero/src/L0Device.cpp
@@ -285,43 +285,45 @@ Error L0DeviceTy::synchronizeImpl(__tgt_async_info &AsyncInfo,
auto &Plugin = getPlugin();
AsyncQueueTy *AsyncQueue = reinterpret_cast<AsyncQueueTy *>(AsyncInfo.Queue);
-
- if (!AsyncQueue->WaitEvents.empty()) {
- const auto &WaitEvents = AsyncQueue->WaitEvents;
- if (Plugin.getOptions().CommandMode == CommandModeTy::AsyncOrdered) {
- // Only need to wait for the last event.
- CALL_ZE_RET_ERROR(zeEventHostSynchronize, WaitEvents.back(),
- L0DefaultTimeout);
- // Synchronize on kernel event to support printf().
- auto KE = AsyncQueue->KernelEvent;
- if (KE && KE != WaitEvents.back()) {
- CALL_ZE_RET_ERROR(zeEventHostSynchronize, KE, L0DefaultTimeout);
- }
- for (auto &Event : WaitEvents) {
- if (auto Err = releaseEvent(Event))
- return Err;
- }
- } else {
- // Async case.
- // Wait for all events. We should wait and reset events in reverse order
- // to avoid premature event reset. If we have a kernel event in the
- // queue, it is the last event to wait for since all wait events of the
- // kernel are signaled before the kernel is invoked. We always invoke
- // synchronization on kernel event to support printf().
- bool WaitDone = false;
- for (auto Itr = WaitEvents.rbegin(); Itr != WaitEvents.rend(); Itr++) {
- if (!WaitDone) {
- CALL_ZE_RET_ERROR(zeEventHostSynchronize, *Itr, L0DefaultTimeout);
- if (*Itr == AsyncQueue->KernelEvent)
- WaitDone = true;
+ {
+ L0DeviceQueueGuard Guard(AsyncQueue, this);
+ if (!AsyncQueue->WaitEvents.empty()) {
+ const auto &WaitEvents = AsyncQueue->WaitEvents;
+ if (Plugin.getOptions().CommandMode == CommandModeTy::AsyncOrdered) {
+ // Only need to wait for the last event.
+ CALL_ZE_RET_ERROR(zeEventHostSynchronize, WaitEvents.back(),
+ L0DefaultTimeout);
+ // Synchronize on kernel event to support printf().
+ auto KE = AsyncQueue->KernelEvent;
+ if (KE && KE != WaitEvents.back()) {
+ CALL_ZE_RET_ERROR(zeEventHostSynchronize, KE, L0DefaultTimeout);
+ }
+ for (auto &Event : WaitEvents) {
+ if (auto Err = releaseEvent(Event))
+ return Err;
+ }
+ } else {
+ // Async case.
+ // Wait for all events. We should wait and reset events in reverse order
+ // to avoid premature event reset. If we have a kernel event in the
+ // queue, it is the last event to wait for since all wait events of the
+ // kernel are signaled before the kernel is invoked. We always invoke
+ // synchronization on kernel event to support printf().
+ bool WaitDone = false;
+ for (auto Itr = WaitEvents.rbegin(); Itr != WaitEvents.rend(); Itr++) {
+ if (!WaitDone) {
+ CALL_ZE_RET_ERROR(zeEventHostSynchronize, *Itr, L0DefaultTimeout);
+ if (*Itr == AsyncQueue->KernelEvent)
+ WaitDone = true;
+ }
+ if (auto Err = releaseEvent(*Itr))
+ return Err;
}
- if (auto Err = releaseEvent(*Itr))
- return Err;
}
+ // In either case, all the events are now reset and released
+ // back into the pool. We need to clear them from the queue.
+ AsyncQueue->WaitEvents.clear();
}
- // In either case, all the events are now reset and released
- // back into the pool. We need to clear them from the queue.
- AsyncQueue->WaitEvents.clear();
}
// Commit delayed USM2M copies.
diff --git a/offload/plugins-nextgen/level_zero/src/L0Memory.cpp b/offload/plugins-nextgen/level_zero/src/L0Memory.cpp
index 477fb0e27c6fc..168fa8b03568a 100644
--- a/offload/plugins-nextgen/level_zero/src/L0Memory.cpp
+++ b/offload/plugins-nextgen/level_zero/src/L0Memory.cpp
@@ -725,8 +725,10 @@ Expected<ze_event_handle_t> EventPoolTy::getEvent() {
/// Return an event to the pool.
Error EventPoolTy::releaseEvent(ze_event_handle_t Event, L0DeviceTy &Device) {
std::lock_guard<std::mutex> Lock(*Mtx);
- CALL_ZE_RET_ERROR(zeEventHostReset, Event);
- Events.push_back(Event);
+ if (std::find(Events.begin(), Events.end(), Event) == Events.end()) {
+ CALL_ZE_RET_ERROR(zeEventHostReset, Event);
+ Events.push_back(Event);
+ }
return Plugin::success();
}
diff --git a/offload/test/api/assert.c b/offload/test/api/assert.c
index 235e2c32c15b7..9a9e770f19622 100644
--- a/offload/test/api/assert.c
+++ b/offload/test/api/assert.c
@@ -1,7 +1,5 @@
// RUN: %libomptarget-compile-run-and-check-generic
// RUN: %libomptarget-compileopt-run-and-check-generic
-// https://github.com/llvm/llvm-project/issues/182119
-// UNSUPPORTED: intelgpu
#include <assert.h>
#include <stdio.h>
diff --git a/offload/test/api/omp_device_managed_memory_alloc.c b/offload/test/api/omp_device_managed_memory_alloc.c
index 098077bdec127..fcacd4015ca49 100644
--- a/offload/test/api/omp_device_managed_memory_alloc.c
+++ b/offload/test/api/omp_device_managed_memory_alloc.c
@@ -1,7 +1,6 @@
// RUN: %libomptarget-compile-run-and-check-generic
// RUN: %libomptarget-compileopt-run-and-check-generic
-// https://github.com/llvm/llvm-project/issues/182119
-// UNSUPPORTED: intelgpu
+// XFAIL: intelgpu
#include <omp.h>
#include <stdio.h>
diff --git a/offload/test/api/omp_host_call.c b/offload/test/api/omp_host_call.c
index b1dd666ea831e..19f4f666cf98a 100644
--- a/offload/test/api/omp_host_call.c
+++ b/offload/test/api/omp_host_call.c
@@ -1,6 +1,5 @@
// RUN: %libomptarget-compile-run-and-check-generic
-// https://github.com/llvm/llvm-project/issues/182119
-// UNSUPPORTED: intelgpu
+// XFAIL: intelgpu
#include <assert.h>
#include <omp.h>
diff --git a/offload/test/api/omp_indirect_func_array.c b/offload/test/api/omp_indirect_func_array.c
index 1064051d412fd..ba329ee034745 100644
--- a/offload/test/api/omp_indirect_func_array.c
+++ b/offload/test/api/omp_indirect_func_array.c
@@ -1,5 +1,5 @@
// RUN: %libomptarget-compile-run-and-check-generic
-// UNSUPPORTED: intelgpu
+// XFAIL: intelgpu
// REQUIRES: gpu
#include <assert.h>
diff --git a/offload/test/api/omp_indirect_func_basic.c b/offload/test/api/omp_indirect_func_basic.c
index b2e29b92f13b3..cbf59ee9172e6 100644
--- a/offload/test/api/omp_indirect_func_basic.c
+++ b/offload/test/api/omp_indirect_func_basic.c
@@ -1,5 +1,5 @@
// RUN: %libomptarget-compile-run-and-check-generic
-// UNSUPPORTED: intelgpu
+// XFAIL: intelgpu
// REQUIRES: gpu
#include <assert.h>
diff --git a/offload/test/api/omp_indirect_func_struct.c b/offload/test/api/omp_indirect_func_struct.c
index 527c9e0b48c5f..2f885577f7c61 100644
--- a/offload/test/api/omp_indirect_func_struct.c
+++ b/offload/test/api/omp_indirect_func_struct.c
@@ -1,5 +1,5 @@
// RUN: %libomptarget-compile-run-and-check-generic
-// UNSUPPORTED: intelgpu
+// XFAIL: intelgpu
// REQUIRES: gpu
#include <assert.h>
diff --git a/offload/test/api/omp_virtual_func.cpp b/offload/test/api/omp_virtual_func.cpp
index 714ba33a513df..dadafbf6907a3 100644
--- a/offload/test/api/omp_virtual_func.cpp
+++ b/offload/test/api/omp_virtual_func.cpp
@@ -1,5 +1,5 @@
// RUN: %libomptarget-compilexx-run-and-check-generic
-// UNSUPPORTED: intelgpu
+// XFAIL: intelgpu
// REQUIRES: gpu
#include <assert.h>
diff --git a/offload/test/api/omp_virtual_func_multiple_inheritance_01.cpp b/offload/test/api/omp_virtual_func_multiple_inheritance_01.cpp
index 373d6fd2af06b..f1e438ce84ecb 100644
--- a/offload/test/api/omp_virtual_func_multiple_inheritance_01.cpp
+++ b/offload/test/api/omp_virtual_func_multiple_inheritance_01.cpp
@@ -1,5 +1,5 @@
// RUN: %libomptarget-compilexx-run-and-check-generic
-// UNSUPPORTED: intelgpu
+// XFAIL: intelgpu
// REQUIRES: gpu
#include <assert.h>
diff --git a/offload/test/api/omp_virtual_func_multiple_inheritance_02.cpp b/offload/test/api/omp_virtual_func_multiple_inheritance_02.cpp
index ae61bfc02d060..9fabd44bb5abe 100644
--- a/offload/test/api/omp_virtual_func_multiple_inheritance_02.cpp
+++ b/offload/test/api/omp_virtual_func_multiple_inheritance_02.cpp
@@ -1,5 +1,5 @@
// RUN: %libomptarget-compilexx-run-and-check-generic
-// UNSUPPORTED: intelgpu
+// XFAIL: intelgpu
// REQUIRES: gpu
#include <assert.h>
diff --git a/offload/test/api/omp_virtual_func_reference.cpp b/offload/test/api/omp_virtual_func_reference.cpp
index deb54b291438c..91e0da38cbcee 100644
--- a/offload/test/api/omp_virtual_func_reference.cpp
+++ b/offload/test/api/omp_virtual_func_reference.cpp
@@ -1,5 +1,5 @@
// RUN: %libomptarget-compilexx-run-and-check-generic
-// UNSUPPORTED: intelgpu
+// XFAIL: intelgpu
// REQUIRES: gpu
#include <assert.h>
diff --git a/offload/test/libc/malloc_parallel.c b/offload/test/libc/malloc_parallel.c
index c74fc6720aadf..42f776826ba08 100644
--- a/offload/test/libc/malloc_parallel.c
+++ b/offload/test/libc/malloc_parallel.c
@@ -1,7 +1,6 @@
// RUN: %libomptarget-compile-run-and-check-generic
// RUN: %libomptarget-compileopt-run-and-check-generic
-// https://github.com/llvm/llvm-project/issues/182119
-// UNSUPPORTED: intelgpu
+// XFAIL: intelgpu
#include <omp.h>
#include <stdio.h>
diff --git a/offload/test/mapping/lambda_by_value.cpp b/offload/test/mapping/lambda_by_value.cpp
index 2978d66c17903..b4e3f16c295a5 100644
--- a/offload/test/mapping/lambda_by_value.cpp
+++ b/offload/test/mapping/lambda_by_value.cpp
@@ -1,7 +1,6 @@
// RUN: %libomptarget-compileopt-generic -fno-exceptions
// RUN: %libomptarget-run-generic 2>&1 | %fcheck-generic
-// https://github.com/llvm/llvm-project/issues/182119
-// UNSUPPORTED: intelgpu
+// XFAIL: intelgpu
#include <stdint.h>
#include <stdio.h>
diff --git a/offload/test/mapping/reduction_implicit_map.cpp b/offload/test/mapping/reduction_implicit_map.cpp
index 7c49f470d7789..d621be142259d 100644
--- a/offload/test/mapping/reduction_implicit_map.cpp
+++ b/offload/test/mapping/reduction_implicit_map.cpp
@@ -1,6 +1,5 @@
// RUN: %libomptarget-compilexx-run-and-check-generic
-// https://github.com/llvm/llvm-project/issues/182119
-// UNSUPPORTED: intelgpu
+// XFAIL: intelgpu
#include <stdio.h>
diff --git a/offload/test/mapping/use_device_addr/target_use_device_addr.c b/offload/test/mapping/use_device_addr/target_use_device_addr.c
index 09bf5a63f3854..00a323f0499db 100644
--- a/offload/test/mapping/use_device_addr/target_use_device_addr.c
+++ b/offload/test/mapping/use_device_addr/target_use_device_addr.c
@@ -1,8 +1,7 @@
// RUN: %libomptarget-compile-generic -fopenmp-version=51
// RUN: %libomptarget-run-generic 2>&1 \
// RUN: | %fcheck-generic
-// https://github.com/llvm/llvm-project/issues/182119
-// UNSUPPORTED: intelgpu
+// XFAIL: intelgpu
#include <stdio.h>
int main() {
diff --git a/offload/test/offloading/assert.cpp b/offload/test/offloading/assert.cpp
index 20d0292d0e399..6ae33d1e6e6cf 100644
--- a/offload/test/offloading/assert.cpp
+++ b/offload/test/offloading/assert.cpp
@@ -1,7 +1,6 @@
// RUN: %libomptarget-compilexx-generic && %libomptarget-run-fail-generic
// RUN: %libomptarget-compileoptxx-generic && %libomptarget-run-fail-generic
-// https://github.com/llvm/llvm-project/issues/182119
-// UNSUPPORTED: intelgpu
+// XFAIL: intelgpu
int main(int argc, char *argv[]) {
#pragma omp target
diff --git a/offload/test/offloading/back2back_distribute.c b/offload/test/offloading/back2back_distribute.c
index d5ab0500fe805..4bd3d52a8c2fb 100644
--- a/offload/test/offloading/back2back_distribute.c
+++ b/offload/test/offloading/back2back_distribute.c
@@ -1,7 +1,6 @@
// clang-format off
// RUN: %libomptarget-compile-generic -O3 && %libomptarget-run-generic | %fcheck-generic
-// https://github.com/llvm/llvm-project/issues/182119
-// UNSUPPORTED: intelgpu
+// XFAIL: intelgpu
// clang-format on
#include <omp.h>
#include <stdio.h>
diff --git a/offload/test/offloading/bug49334.cpp b/offload/test/offloading/bug49334.cpp
index c7554b109c431..47b79d408b93b 100644
--- a/offload/test/offloading/bug49334.cpp
+++ b/offload/test/offloading/bug49334.cpp
@@ -8,8 +8,7 @@
// REQUIRES: gpu
// UNSUPPORTED: nvidiagpu
// UNSUPPORTED: amdgpu
-// https://github.com/llvm/llvm-project/issues/182119
-// UNSUPPORTED: intelgpu
+// XFAIL: intelgpu
#include <cassert>
#include <cmath>
diff --git a/offload/test/offloading/bug51982.c b/offload/test/offloading/bug51982.c
index ca32ff29437c9..a6910b1fbb3f8 100644
--- a/offload/test/offloading/bug51982.c
+++ b/offload/test/offloading/bug51982.c
@@ -1,8 +1,7 @@
// RUN: %libomptarget-compile-generic -O2 && %libomptarget-run-generic
// -O2 to run openmp-opt
// RUN: %libomptarget-compileopt-generic -O2 && %libomptarget-run-generic
-// https://github.com/llvm/llvm-project/issues/182119
-// UNSUPPORTED: intelgpu
+// XFAIL: intelgpu
int main(void) {
long int aa = 0;
diff --git a/offload/test/offloading/bug74582.c b/offload/test/offloading/bug74582.c
index 247f91ee5183f..c6a283bb93691 100644
--- a/offload/test/offloading/bug74582.c
+++ b/offload/test/offloading/bug74582.c
@@ -1,7 +1,5 @@
// RUN: %libomptarget-compile-generic && %libomptarget-run-generic
// RUN: %libomptarget-compileopt-generic && %libomptarget-run-generic
-// https://github.com/llvm/llvm-project/issues/182119
-// UNSUPPORTED: intelgpu
// Verify we do not read bits in the image that are not there (nobits section).
diff --git a/offload/test/offloading/ctor_dtor_api.cpp b/offload/test/offloading/ctor_dtor_api.cpp
index fb42f646654c4..cbb78cff947af 100644
--- a/offload/test/offloading/ctor_dtor_api.cpp
+++ b/offload/test/offloading/ctor_dtor_api.cpp
@@ -1,7 +1,6 @@
// RUN: %libomptarget-compilexx-run-and-check-generic
// RUN: %libomptarget-compileoptxx-run-and-check-generic
-// https://github.com/llvm/llvm-project/issues/182119
-// UNSUPPORTED: intelgpu
+// XFAIL: intelgpu
#include <cstdio>
#include <omp.h>
diff --git a/offload/test/offloading/ctor_dtor_lazy.cpp b/offload/test/offloading/ctor_dtor_lazy.cpp
index 015d8877e75e4..d09bd517387de 100644
--- a/offload/test/offloading/ctor_dtor_lazy.cpp
+++ b/offload/test/offloading/ctor_dtor_lazy.cpp
@@ -4,8 +4,7 @@
// RUN: %not --crash %libomptarget-run-generic
// RUN: %libomptarget-compilexx-generic -DCTOR_API
// RUN: %not --crash %libomptarget-run-generic
-// https://github.com/llvm/llvm-project/issues/182119
-// UNSUPPORTED: intelgpu
+// XFAIL: intelgpu
#include <cstdio>
#include <omp.h>
diff --git a/offload/test/offloading/dyn_groupprivate.cpp b/offload/test/offloading/dyn_groupprivate.cpp
index fd0c3de0c8c5d..5a478a0c5a917 100644
--- a/offload/test/offloading/dyn_groupprivate.cpp
+++ b/offload/test/offloading/dyn_groupprivate.cpp
@@ -3,7 +3,7 @@
// RUN: %libomptarget-compileoptxx-generic -fopenmp-version=61
// RUN: %libomptarget-run-generic | %fcheck-generic
// REQUIRES: gpu
-// UNSUPPORTED: intelgpu
+// XFAIL: intelgpu
#include <omp.h>
#include <stdio.h>
diff --git a/offload/test/offloading/dynamic_module.c b/offload/test/offloading/dynamic_module.c
index f3f59743d3f0e..3df81eb03648a 100644
--- a/offload/test/offloading/dynamic_module.c
+++ b/offload/test/offloading/dynamic_module.c
@@ -5,8 +5,7 @@
// RUN: %libomptarget-compileopt-generic %t.so && %libomptarget-run-generic 2>&1 | %fcheck-generic
//
// REQUIRES: gpu
-// https://github.com/llvm/llvm-project/issues/182119
-// UNSUPPORTED: intelgpu
+
// clang-format on
#ifdef SHARED
diff --git a/offload/test/offloading/high_trip_count_block_limit.cpp b/offload/test/offloading/high_trip_count_block_limit.cpp
index 165ff79f3482e..ed6fd8fa8b33e 100644
--- a/offload/test/offloading/high_trip_count_block_limit.cpp
+++ b/offload/test/offloading/high_trip_count_block_limit.cpp
@@ -5,8 +5,7 @@
// UNSUPPORTED: aarch64-unknown-linux-gnu
// UNSUPPORTED: x86_64-unknown-linux-gnu
// UNSUPPORTED: s390x-ibm-linux-gnu
-// https://github.com/llvm/llvm-project/issues/182119
-// UNSUPPORTED: intelgpu
+
// clang-format on
/*
diff --git a/offload/test/offloading/malloc.c b/offload/test/offloading/malloc.c
index fd912b8188f3a..d9b87d2bbe27c 100644
--- a/offload/test/offloading/malloc.c
+++ b/offload/test/offloading/malloc.c
@@ -1,7 +1,6 @@
// RUN: %libomptarget-compile-generic && %libomptarget-run-generic
// RUN: %libomptarget-compileopt-generic && %libomptarget-run-generic
-// https://github.com/llvm/llvm-project/issues/182119
-// UNSUPPORTED: intelgpu
+// XFAIL: intelgpu
#include <stdio.h>
#include <stdlib.h>
diff --git a/offload/test/offloading/multiple_reductions_simple.c b/offload/test/offloading/multiple_reductions_simple.c
index 87efb3f47dd53..5357caa46c29b 100644
--- a/offload/test/offloading/multiple_reductions_simple.c
+++ b/offload/test/offloading/multiple_reductions_simple.c
@@ -1,7 +1,6 @@
// RUN: %libomptarget-compile-run-and-check-generic
// RUN: %libomptarget-compileopt-run-and-check-generic
-// https://github.com/llvm/llvm-project/issues/182119
-// UNSUPPORTED: intelgpu
+// XFAIL: intelgpu
#include <stdio.h>
diff --git a/offload/test/offloading/ompx_coords.c b/offload/test/offloading/ompx_coords.c
index b40ae572a7704..71979ef2ebe87 100644
--- a/offload/test/offloading/ompx_coords.c
+++ b/offload/test/offloading/ompx_coords.c
@@ -1,8 +1,7 @@
// RUN: %libomptarget-compileopt-run-and-check-generic
//
// REQUIRES: gpu
-// https://github.com/llvm/llvm-project/issues/182119
-// UNSUPPORTED: intelgpu
+// XFAIL: intelgpu
#include <omp.h>
#include <ompx.h>
diff --git a/offload/test/offloading/ompx_saxpy_mixed.c b/offload/test/offloading/ompx_saxpy_mixed.c
index 0386c568b23f0..f370d68327c93 100644
--- a/offload/test/offloading/ompx_saxpy_mixed.c
+++ b/offload/test/offloading/ompx_saxpy_mixed.c
@@ -1,8 +1,7 @@
// RUN: %libomptarget-compileopt-run-and-check-generic
//
// REQUIRES: gpu
-// https://github.com/llvm/llvm-project/issues/182119
-// UNSUPPORTED: intelgpu
+// XFAIL: intelgpu
#include <math.h>
#include <omp.h>
diff --git a/offload/test/offloading/parallel_target_teams_reduction.cpp b/offload/test/offloading/parallel_target_teams_reduction.cpp
index 335cee994b19c..05f675b87c5e1 100644
--- a/offload/test/offloading/parallel_target_teams_reduction.cpp
+++ b/offload/test/offloading/parallel_target_teams_reduction.cpp
@@ -3,8 +3,7 @@
// FIXME: This is a bug in host offload, this should run fine.
// REQUIRES: gpu
-// https://github.com/llvm/llvm-project/issues/182119
-// UNSUPPORTED: intelgpu
+// XFAIL: intelgpu
#include <iostream>
#include <vector>
diff --git a/offload/test/offloading/parallel_target_teams_reduction_max.cpp b/offload/test/offloading/parallel_target_teams_reduction_max.cpp
index 4654b77188641..bb893e6f3f755 100644
--- a/offload/test/offloading/parallel_target_teams_reduction_max.cpp
+++ b/offload/test/offloading/parallel_target_teams_reduction_max.cpp
@@ -3,8 +3,7 @@
// FIXME: This is a bug in host offload, this should run fine.
// REQUIRES: gpu
-// https://github.com/llvm/llvm-project/issues/182119
-// UNSUPPORTED: intelgpu
+// XFAIL: intelgpu
// This test validates that the OpenMP target reductions to find a maximum work
// as intended for a few common data types.
diff --git a/offload/test/offloading/parallel_target_teams_reduction_min.cpp b/offload/test/offloading/parallel_target_teams_reduction_min.cpp
index 0172b440e0648..416ef44f21127 100644
--- a/offload/test/offload...
[truncated]
|
kevinsala
reviewed
Mar 17, 2026
| : queue(Queue), device(Device) {} | ||
| ~L0DeviceQueueGuard() { | ||
| for (auto &Event : queue->WaitEvents) | ||
| auto Err = device->releaseEvent(Event); |
Contributor
There was a problem hiding this comment.
This needs proper error propagation. An output Error (pointer or reference) within the object could do the trick. Similar to AsyncInfoWrapperTy.
Contributor
Author
There was a problem hiding this comment.
The whole point of this is to do the clean up in case synchronizeImpl function encountered error and ignore errors during the clean up as the original error is reported by synchronizeImpl. In normal circumstances this code won't be executed as WaitEvents list is empty.
I am not sure where to propagate errors from here as it is a destructor of a stack object inside a function. Nobody calls it and nobody expects an error from this object
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
3 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
This change improves handling of errors during synchronization in Level Zero plugin by ensuring cleanup of queues and events in case of an synchronization error. As a result multiple tests stopped hanging.