[SYCL][Graph][UR][Level Zero] Executable-graph finalization causes unbounded cached zeEvent retention on Windows
- Dominant language
- LLVM
- Stars
- 1.5k
- Forks
- 854
- Avg merge
- 3d 17h
- Merged PRs (30d)
- 137
Description
### Describe the bug
On an Intel Arc A770 under the Windows / oneAPI / graphics-driver stack listed
below, repeatedly **constructing, finalizing, and destroying** non-updatable
experimental SYCL executable graphs causes approximately linear host-process
private-memory growth when Unified Runtime Level Zero event caching is enabled
(the default).
The retained resource has been localized to **Level Zero events held by the
Unified Runtime Level Zero event-cache/recycling path**:
- plain kernel submission is memory-flat;
- graph recording and destruction without finalization is memory-flat;
- finalization followed by destruction grows by approximately **1.5 KiB per
captured node per graph generation**;
- submitting the executable graph does not add to the finalization-only slope;
- `UR_L0_LEAKS_DEBUG=1` shows approximately one additional unmatched
`zeEvent` per captured node per finalized graph;
- `UR_L0_DISABLE_EVENTS_CACHING=1` balances event create/destroy calls and
makes host private memory flat; and
- `SYCL_UR_TRACE=2` shows that `urCommandBufferReleaseExp` is called after the
SYCL executable graph is destroyed.
The evidence therefore excludes a missing command-buffer release in the SYCL
graph destructor. The observable defect is below that API boundary: events
created for the finalized command buffer enter the UR Level Zero event cache,
but repeated equivalent graph generations do not reuse or drain them in a
bounded way.
This package is self-contained. It depends only on the Intel oneAPI DPC++
toolchain, `` plus the experimental graph header, the Windows
`GetProcessMemoryInfo` API, and a single durable USM allocation. It has no
model, ggml, application, private DLL, or third-party source dependency.
> The measured value is **host process private memory** (committed private
> bytes), **not GPU-local VRAM**.
### To reproduce
A complete hardened reproducer, including four lifecycle modes, Windows build/run scripts, reference measurements, the event-cache A/B, and the Unified Runtime trace, is available here:
**Full reproducer repo:** [https://github.com/congdm/intel-sycl-graph-finalize-retention-issue-report/tree/343a623f35aa2d815a4f373c3d43562b2cb0fb41](https://github.com/congdm/intel-sycl-graph-finalize-retention-issue-report/tree/343a623f35aa2d815a4f373c3d43562b2cb0fb41)
The following is the minimized finalization-only case. The executable graph is never submitted; it is finalized and destroyed before each PrivateUsage sample.
```cpp
#include
#include
#include
#include
#include
namespace sycl_exp = sycl::ext::oneapi::experimental;
static double private_mib() {
PROCESS_MEMORY_COUNTERS_EX pmc = {};
GetProcessMemoryInfo(
GetCurrentProcess(),
reinterpret_cast(&pmc),
sizeof(pmc));
return static_cast(pmc.PrivateUsage) / (1024.0 * 1024.0);
}
int main() {
constexpr int iterations = 100;
constexpr int nodes = 1024;
sycl::queue queue{sycl::default_selector_v};
int * output = sycl::malloc_shared(nodes, queue);
// Compile/warm the kernel before the measured graph-finalization loop.
queue.parallel_for(sycl::range<1>(nodes), [=](sycl::id<1> id) {
output[id[0]] = 0;
}).wait_and_throw();
std::printf("iteration private_MiB\n");
std::printf("0 %.3f\n", private_mib());
for (int iteration = 1; iteration <= iterations; ++iteration) {
{
sycl_exp::command_graph graph(
queue,
{sycl_exp::property::graph::assume_buffer_outlives_graph{}});
graph.begin_recording(queue);
for (int node = 0; node < nodes; ++node) {
queue.single_task([=]() {
output[node] = node;
});
}
graph.end_recording();
// The executable is finalized but never submitted.
auto executable = graph.finalize();
}
// Both the executable and modifiable graph have been destroyed.
std::printf("%d %.3f\n", iteration, private_mib());
}
sycl::free(output, queue);
}
```
### Compile
From a oneAPI 2026.1 command environment on Windows:
```bat
"C:\Program Files (x86)\Intel\oneAPI\compiler\2026.1\bin\icx.exe" -fsycl /EHsc -O2 -DNDEBUG /std:c++17 repro.cpp -o repro.exe
```
### Run with the default Unified Runtime Level Zero event cache
```bat
repro.exe
```
### Actual behavior
On Intel Arc A770, the process PrivateUsage grows approximately linearly by about 1.5 MiB per iteration for a 1024-node graph, or approximately 1.5 KiB per finalized graph node.
The complete reproducer additionally shows:
- direct kernel submission: memory-flat;
- graph record/destroy without finalization: memory-flat;
- finalize/destroy without submission: approximately +1.5 MiB/iteration;
- finalize/submit/destroy: approximately the same slope.
Submission is therefore not required.
### Event-cache control
Run the same executable with event caching disabled:
```bat
set UR_L0_DISABLE_EVENTS_CACHING=1
repro.exe
```
### Expected behavior
Destroying the executable graph should release its resources or return them to a bounded/reusable cache. Repeated destruction of structurally identical finalized graphs should not cause unbounded process-private memory growth.
### Actual behavior with event caching disabled
PrivateUsage becomes flat.
With `UR_L0_LEAKS_DEBUG=1`, the 64-node, five-iteration hardened reproducer reports:
```text
Default cache, finalize_destroy:
zeEventCreate = 382
zeEventDestroy = 0
Event cache disabled, finalize_destroy:
zeEventCreate = 591
zeEventDestroy = 591
```
The direct and record-only controls are memory-flat. The mode-dependent difference is approximately one additional cached `zeEvent` per finalized graph node per graph generation.
### Unified Runtime ownership trace
For a four-node, one-finalization run with:
```bat
set SYCL_UR_TRACE=2
```
the trace contains:
```text
urCommandBufferCreateExp 1
urCommandBufferAppendKernelLaunchExp 4
urCommandBufferFinalizeExp 1
urCommandBufferReleaseExp 1
```
`urCommandBufferReleaseExp` is present after destruction of the SYCL executable graph. This excludes a missing command-buffer release at the SYCL-to-Unified-Runtime API boundary.
Combined with the cache-disabled result, the observable defect is in the Unified Runtime Level Zero adapter’s event-cache/recycling behavior around a released executable command buffer.
The behavior reproduces with immediate command lists both enabled and disabled and does not require a host task or executable-graph submission.
### Environment
- **OS:** Microsoft Windows 10 IoT Enterprise LTSC 2021, version 21H2, build 19044.7663
- **Target device and vendor:** Intel(R) Arc(TM) A770 Graphics, Intel(R) Corporation; `intel_gpu_acm_g10`; Unified Runtime Level Zero backend
- **Windows graphics driver:** 32.0.101.8974
- **DPC++ version:**
```text
Intel(R) oneAPI DPC++/C++ Compiler for applications running on Intel(R) 64,
Version 2026.1.0 Build 20260617
```
- **Dependencies / `sycl-ls --verbose`:**
```text
Platform:
Version : 1.15
Name : Intel(R) oneAPI Unified Runtime over Level-Zero
Vendor : Intel(R) Corporation
Device:
Type : gpu
Version : 12.55.8
Name : Intel(R) Arc(TM) A770 Graphics
Vendor : Intel(R) Corporation
Driver : 1.15.39183+1
DeviceID : 22176
Architecture : intel_gpu_acm_g10
Graph aspects:
ext_oneapi_limited_graph = true
ext_oneapi_graph = false
Default selector:
[level_zero:gpu:0] Intel(R) oneAPI Unified Runtime over Level-Zero,
Intel(R) Arc(TM) A770 Graphics 12.55.8 [1.15.39183+1]
Loaded Unified Runtime adapters:
ur_adapter_level_zero.dll
ur_adapter_level_zero_v2.dll
ur_adapter_opencl.dll
```
### Additional context
_No response_
Contributor guide
Assessment
This issue has not been assessed yet.