Repository navigation
[SYCL] Destructors in work-group-scope causing issues. #23461
Description
Describe the bug
When using parallel_for_work_group with a lambda capturing a non-trivial destructor, the compiled program simply hangs. I believe this is due to incorrect assumptions over the availability of __cxa_atexit in this context.
To reproduce
Using the following code:
#include <sycl/sycl.hpp>
#include <iostream>
struct Counter {
int *Ptr;
~Counter() { *Ptr += 1; }
};
int main() {
sycl::queue Q;
int *Result = sycl::malloc_shared<int>(1, Q);
*Result = 0;
Q.submit([&](sycl::handler &CGH) {
CGH.parallel_for_work_group(sycl::range<1>{1}, sycl::range<1>{1},
[=](sycl::group<1>) {
// Work-group scope variable with a
// non-trivial destructor.
Counter C{Result};
});
}).wait();
int Value = *Result;
sycl::free(Result, Q);
std::cout << "Result = " << Value << " (expected 1)\n";
return Value == 1 ? 0 : 1;
}
Running clang -fsycl -Wall -Wextra -Wpedantic [CODE] && ./a.out causes the executable to hang.
Running clang -fsycl -fsycl-device-only -S -emit-llvm [CODE] shows that the device side calls __cxa_atexit, which I don't believe is supported in this context.
I believe compiling this code should produce at either a warning or an error.
Environment
- OS: Linux
- Target device and vendor: Intel CPU
- DPC++ version: DPC++ compiler 7.2.0 (pre-release) build based on:
- clang version 24.0.0git (https://github.com/inteltdelisle/llvm.git e35fca9)
- Target: x86_64-unknown-linux-gnu
- Thread model: posix
- Build config: +assertions
Additional context
This was found while working on #23250. That PR changes the behavior of this bug to silently drop the destructor, which still should be behind a warning or error.
Activity
gmlueck commentedon Oct 13, 2026
This is an interesting case, and I do not think the SYCL 2020 spec is clear about the intended semantic. However, it certainly does not make sense to call
__cxa_atexitin this case. I believe that function is used to register a destructor that runs at program exit, right? Even if there was some notion of "program exit" for a kernel, it would not make sense to destroy work-group local variables this way.The lifetime of a work-group local variable logically ends when the last work-item in the work-group completes. Therefore, if we were going to run the destructor, we would have to somehow call it when the last work-item in the work-group completes. Correspondingly, we would need to run the constructor when the first work-item in the work-group starts executing.
However, this is not how objects in a
local_accessorbehave. The SYCL 2020 spec is clear about this case. If the type of the object in a local accessor is not an "implicit-lifetime type", the local accessor merely allocates storage for the object. In this case, it is the user's responsibility to call the constructor / destructor manually.If we were to treat work-group local objects in hierarchical parallelism the same, then the
Cvariable in your example above would silently skip calling its destructor (and would silently skip the constructor if it had one). However, that seems like it would be pretty confusing.Instead, I think it would make sense to make it illegal to declare a work-group local variable whose type has a non-trivial constructor or destructor.