[SYCL][Graph][UR][Level Zero] Executable-graph finalization causes unbounded cached zeEvent retention on Windows

[congdm]

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, <sycl/sycl.hpp> 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

The following is the minimized finalization-only case. The executable graph is never submitted; it is finalized and destroyed before each PrivateUsage sample.

#include <sycl/sycl.hpp>
#include <sycl/ext/oneapi/experimental/graph/command_graph.hpp>

#include <windows.h>
#include <psapi.h>
#include <cstdio>

namespace sycl_exp = sycl::ext::oneapi::experimental;

static double private_mib() {
    PROCESS_MEMORY_COUNTERS_EX pmc = {};
    GetProcessMemoryInfo(
        GetCurrentProcess(),
        reinterpret_cast<PROCESS_MEMORY_COUNTERS *>(&pmc),
        sizeof(pmc));
    return static_cast<double>(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<int>(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:

"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

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:

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:

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:

set SYCL_UR_TRACE=2

the trace contains:

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:

Intel(R) oneAPI DPC++/C++ Compiler for applications running on Intel(R) 64,
Version 2026.1.0 Build 20260617

• Dependencies / sycl-ls --verbose:

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

[bratpiorka]

hi @congdm could you check if the issue still occurs when using a newer version of the Level Zero adapter?
It can be enabled on Windows with (see https://intel.github.io/llvm/EnvironmentVariables.html#controlling-dpc-level-zero-adapter)

[congdm]

hi @congdm could you check if the issue still occurs when using a newer version of the Level Zero adapter? It can be enabled on Windows with (see https://intel.github.io/llvm/EnvironmentVariables.html#controlling-dpc-level-zero-adapter)

Oh, after the flag SYCL_UR_USE_LEVEL_ZERO_V2 is set, the memory growth indeed is flat instead of leaking through time. But from the documentation, I see that this feature is intended for Battlemage onward, is there side-effect for using it with A770?

Results:

Mode	Default slope	V2 slope	Default first→last MiB	V2 first→last MiB
direct	0.0000	0.0000	146.62 → 146.94	151.64 → 151.70
record_destroy	0.0181	0.0062	150.48 → 152.44	151.71 → 152.90
finalize_destroy	1.4862	-0.0007	146.73 → 302.94	151.62 → 162.50
submit_destroy	1.4854	0.2553	149.77 → 306.09	151.78 → 172.46

Key finding

• finalize_destroy — the primary leak signal — is flat under V2. It climbs ~11 MiB in the first few iterations (warm-up) then stays pinned at ~162.5–164 MiB; second-half slope = -0.0007.
• submit_destroy — the apparent 0.2553 slope is not linear growth. It's a step: flat near ~163 MiB, then a single step to ~172.5 MiB at iter ~70, then flat again. Slopes inside the flat segments are [50..65] = 0.0219 and [70..100] = 0.0224 — noise-level. There is no continuous per-iteration retention.

So the unbounded, iteration-linear Level Zero event-cache retention (≈1.5 MiB/generation) is eliminated by the Level Zero V2 adapter on this lifecycle.