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);
 }