From 43e0d765fcb74644b7dbd31b1f93480d0e9e919d Mon Sep 17 00:00:00 2001 From: Mateusz Hoppe Date: Tue, 20 Oct 2026 14:59:08 +0200 Subject: [PATCH] [SYCL][E2E] Wait for kernels in max_num_work_groups_sync_query test testFreeFunctionKernel() launches a kernel with use_root_sync and the maximum number of work-groups returned by max_num_work_groups_sync, then frees the USM allocation used by the kernel without waiting for the kernel to finish. The test runs on a default (out-of-order) queue, so the four free-function launches are independent and may execute concurrently. Each launch is sized by max_num_work_groups_sync to occupy the whole device, so when two of them overlap neither can have all of its work-groups resident at once and their root-group barriers never complete. This was observed as a hang on Level Zero, where the out-of-order queue distributes the launches over several command lists and hardware contexts. Freeing memory that a running kernel still writes is also undefined behavior. Wait on the queue before freeing the allocation, the same way testLambdaKernel() already does. This serializes the root-group synchronized launches and removes the use-after-free. HSD-ES: 21045664176 Signed-off-by: Mateusz Hoppe --- .../GroupAlgorithm/max_num_work_groups_sync_query.cpp | 5 +++++ 1 file changed, 5 insertions(+) diff --git a/sycl/test-e2e/GroupAlgorithm/max_num_work_groups_sync_query.cpp b/sycl/test-e2e/GroupAlgorithm/max_num_work_groups_sync_query.cpp index 61905f19ba0b0..0732d8c97a3f0 100644 --- a/sycl/test-e2e/GroupAlgorithm/max_num_work_groups_sync_query.cpp +++ b/sycl/test-e2e/GroupAlgorithm/max_num_work_groups_sync_query.cpp @@ -95,6 +95,11 @@ void testFreeFunctionKernel(sycl::queue &q, sycl::context &ctx, } catch (sycl::exception &e) { assert(maxWGs == 0); } + // Wait for the kernel before freeing its memory. This also keeps + // root-group-synchronized kernels from running concurrently on an + // out-of-order queue: each is sized to fill the device, so overlapping + // launches cannot all have their work-groups resident and deadlock. + q.wait(); sycl::free(data, ctx); }