[SYCL] `get_opportunistic_group()` and `group_ballot()` include returned work-items at `-O0`

### Describe the bug

On an Intel Data Center GPU Max 1550 (Ponte Vecchio), when the device code is compiled **without optimisations** (`-O0`), work-items that have already returned from the kernel are still treated as members of the group:

- `sycl::ext::oneapi::experimental::this_work_item::get_opportunistic_group()` returns a `fragment` whose `get_local_linear_range()` is the size of the whole sub-group, instead of the number of work-items that actually called the function;
- `sycl::ext::oneapi::group_ballot(sub_group, true)` returns a mask with the bits of the work-items that have returned set to 1.

The same code compiled with `-O2` produces the correct results.

According to the [`sycl_ext_oneapi_non_uniform_groups`](https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/experimental/sycl_ext_oneapi_non_uniform_groups.asciidoc) specification, `this_work_item::get_opportunistic_group()`

> _Returns_: A `fragment` consisting of all work-items in the same sub-group as the calling work-item which call this function simultaneously.

so work-items that have returned from the kernel should not be part of the `fragment`.

In DPC++, on SPIR-V targets, `get_opportunistic_group()` is implemented on top of `group_ballot(this_sub_group(), true)` (see `sycl/include/sycl/ext/oneapi/experimental/fragment.hpp`), which is consistent with both functions showing the same wrong behaviour.

The results should not depend on the optimisation level: `get_opportunistic_group()` should return a `fragment` containing only the work-items that call it, and `group_ballot()` should not report as set the bits of the work-items that have returned from the kernel, as it happens with `-O2`.

### To reproduce

1. Save the following code as `test.cc`:

```cpp
#include

#include
#include
#include
#include

// oneAPI 2026.0 replaced this_kernel::get_opportunistic_group(), returning an opportunistic_group, with
// this_work_item::get_opportunistic_group(), returning a fragment (intel/llvm#19238).
// __INTEL_LLVM_COMPILER encodes the oneAPI version as YYYYMMPP, e.g. 20250300 for 2025.3.0 and 20260000 for 2026.0.0.
#if defined(__INTEL_LLVM_COMPILER) && __INTEL_LLVM_COMPILER >= 20260000
using sycl::ext::oneapi::experimental::this_work_item::get_opportunistic_group;
#else
using sycl::ext::oneapi::experimental::this_kernel::get_opportunistic_group;
#endif

// Run a kernel with 2 sub-groups of size SgSize.
// The work items with a local id >= SgSize / 2 return early, the others call group_ballot(sg, true)
// and get_opportunistic_group(), and record the results.
// If `converged` is true, all work items call group_ballot once before diverging.
template
bool run(sycl::queue& q, bool converged)
{
constexpr std::size_t groups = 2;
std::size_t const n = SgSize * groups;
std::uint32_t const expectedMask = (std::uint32_t{1} << (SgSize / 2)) - 1u;
std::uint32_t const expectedCount = SgSize / 2;

auto* mask = sycl::malloc_device(n, q);
auto* count = sycl::malloc_device(n, q);
q.fill(mask, std::uint32_t{0xdeadbeef}, n).wait();
q.fill(count, std::uint32_t{0}, n).wait();

q.parallel_for(
sycl::nd_range<1>{n, SgSize},
[=](sycl::nd_item<1> item) [[sycl::reqd_sub_group_size(SgSize)]]
{
auto sg = item.get_sub_group();
std::size_t const i = item.get_global_linear_id();
std::size_t const lid = sg.get_local_linear_id();

if(converged)
{
[[maybe_unused]] auto all = sycl::ext::oneapi::group_ballot(sg, true);
}

// half of the work items leave the kernel
if(lid >= SgSize / 2)
return;

auto m = sycl::ext::oneapi::group_ballot(sg, true);
std::uint32_t bits = 0;
m.extract_bits(bits);
mask[i] = bits;

auto og = get_opportunistic_group();
count[i] = static_cast(og.get_local_linear_range());
})
.wait();

std::vector masks(n), counts(n);
q.memcpy(masks.data(), mask, n * sizeof(std::uint32_t)).wait();
q.memcpy(counts.data(), count, n * sizeof(std::uint32_t)).wait();
sycl::free(mask, q);
sycl::free(count, q);

bool ok = true;
std::cout << "sub-group size " << SgSize << ", " << (converged ? "converged first" : "diverged") << '\n';
std::cout << " expected: ballot mask 0x" << std::hex << expectedMask << std::dec << ", opportunistic group size "
<< expectedCount << '\n';
for(std::size_t i = 0; i < n; ++i)
{
if(i % SgSize >= SgSize / 2)
continue;
bool const good = masks[i] == expectedMask and counts[i] == expectedCount;
ok = ok and good;
if(not good)
std::cout << " item " << i << " (sub-group " << i / SgSize << ", local id " << i % SgSize
<< "): ballot mask 0x" << std::hex << masks[i] << std::dec << ", opportunistic group size "
<< counts[i] << '\n';
}
std::cout << " " << (ok ? "OK" : "FAILED") << '\n';
return ok;
}

int main()
{
sycl::queue q{sycl::gpu_selector_v};
auto const dev = q.get_device();
std::cout << "Device: " << dev.get_info() << '\n';
auto const sizes = dev.get_info();

bool ok = true;
for(bool converged : {false, true})
{
if(std::find(sizes.begin(), sizes.end(), 16) != sizes.end())
ok = run<16>(q, converged) and ok;
if(std::find(sizes.begin(), sizes.end(), 32) != sizes.end())
ok = run<32>(q, converged) and ok;
}
return ok ? 0 : 1;
}
```

The kernel is run over 2 sub-groups of size 16 and 32. The work-items with a local id `>= SgSize / 2` return immediately; the other work-items call `group_ballot(sg, true)` and `get_opportunistic_group()`, and record the resulting mask and the size of the opportunistic group. The expected results are a mask with the lower `SgSize / 2` bits set, and an opportunistic group of size `SgSize / 2`.

Each configuration is run twice: with the work-items diverging immediately, and with all work-items calling `group_ballot(sg, true)` once in converged control flow before diverging.

2. Compile with optimisations and run:

```
$ icpx -fsycl -fsycl-targets=spir64_gen -Xs "-device pvc" -O2 test.cc -o test_O2
Compilation from IR - skipping loading of FCL
Build succeeded.
Compilation from IR - skipping loading of FCL
Build succeeded.
Compilation from IR - skipping loading of FCL
Build succeeded.

$ ./test_O2
Device: Intel(R) Data Center GPU Max 1550
sub-group size 16, diverged
expected: ballot mask 0xff, opportunistic group size 8
OK
sub-group size 32, diverged
expected: ballot mask 0xffff, opportunistic group size 16
OK
sub-group size 16, converged first
expected: ballot mask 0xff, opportunistic group size 8
OK
sub-group size 32, converged first
expected: ballot mask 0xffff, opportunistic group size 16
OK
```

All the tests pass.

3. Compile without optimisations and run:

```
$ icpx -fsycl -fsycl-targets=spir64_gen -Xs "-device pvc" -O0 test.cc -o test_O0
Compilation from IR - skipping loading of FCL
Build succeeded.
Compilation from IR - skipping loading of FCL
Build succeeded.
Compilation from IR - skipping loading of FCL
Build succeeded.
```

```
$ ./test_O0
Device: Intel(R) Data Center GPU Max 1550
sub-group size 16, diverged
expected: ballot mask 0xff, opportunistic group size 8
item 0 (sub-group 0, local id 0): ballot mask 0xffff, opportunistic group size 16
item 1 (sub-group 0, local id 1): ballot mask 0xffff, opportunistic group size 16
...
item 23 (sub-group 1, local id 7): ballot mask 0xffff, opportunistic group size 16
FAILED
sub-group size 32, diverged
expected: ballot mask 0xffff, opportunistic group size 16
item 0 (sub-group 0, local id 0): ballot mask 0xffffffff, opportunistic group size 32
...
item 47 (sub-group 1, local id 15): ballot mask 0xffffffff, opportunistic group size 32
FAILED
sub-group size 16, converged first
...
FAILED
sub-group size 32, converged first
...
FAILED
```

All the tests fail: every work-item that did not return sees a ballot mask with all the bits set (`0xffff` or `0xffffffff`) and an opportunistic group containing the whole sub-group (16 or 32 work-items).

Full output of ./test_O0

```
$ ./test_O0
Device: Intel(R) Data Center GPU Max 1550
sub-group size 16, diverged
expected: ballot mask 0xff, opportunistic group size 8
item 0 (sub-group 0, local id 0): ballot mask 0xffff, opportunistic group size 16
item 1 (sub-group 0, local id 1): ballot mask 0xffff, opportunistic group size 16
item 2 (sub-group 0, local id 2): ballot mask 0xffff, opportunistic group size 16
item 3 (sub-group 0, local id 3): ballot mask 0xffff, opportunistic group size 16
item 4 (sub-group 0, local id 4): ballot mask 0xffff, opportunistic group size 16
item 5 (sub-group 0, local id 5): ballot mask 0xffff, opportunistic group size 16
item 6 (sub-group 0, local id 6): ballot mask 0xffff, opportunistic group size 16
item 7 (sub-group 0, local id 7): ballot mask 0xffff, opportunistic group size 16
item 16 (sub-group 1, local id 0): ballot mask 0xffff, opportunistic group size 16
item 17 (sub-group 1, local id 1): ballot mask 0xffff, opportunistic group size 16
item 18 (sub-group 1, local id 2): ballot mask 0xffff, opportunistic group size 16
item 19 (sub-group 1, local id 3): ballot mask 0xffff, opportunistic group size 16
item 20 (sub-group 1, local id 4): ballot mask 0xffff, opportunistic group size 16
item 21 (sub-group 1, local id 5): ballot mask 0xffff, opportunistic group size 16
item 22 (sub-group 1, local id 6): ballot mask 0xffff, opportunistic group size 16
item 23 (sub-group 1, local id 7): ballot mask 0xffff, opportunistic group size 16
FAILED
sub-group size 32, diverged
expected: ballot mask 0xffff, opportunistic group size 16
item 0 (sub-group 0, local id 0): ballot mask 0xffffffff, opportunistic group size 32
item 1 (sub-group 0, local id 1): ballot mask 0xffffffff, opportunistic group size 32
item 2 (sub-group 0, local id 2): ballot mask 0xffffffff, opportunistic group size 32
item 3 (sub-group 0, local id 3): ballot mask 0xffffffff, opportunistic group size 32
item 4 (sub-group 0, local id 4): ballot mask 0xffffffff, opportunistic group size 32
item 5 (sub-group 0, local id 5): ballot mask 0xffffffff, opportunistic group size 32
item 6 (sub-group 0, local id 6): ballot mask 0xffffffff, opportunistic group size 32
item 7 (sub-group 0, local id 7): ballot mask 0xffffffff, opportunistic group size 32
item 8 (sub-group 0, local id 8): ballot mask 0xffffffff, opportunistic group size 32
item 9 (sub-group 0, local id 9): ballot mask 0xffffffff, opportunistic group size 32
item 10 (sub-group 0, local id 10): ballot mask 0xffffffff, opportunistic group size 32
item 11 (sub-group 0, local id 11): ballot mask 0xffffffff, opportunistic group size 32
item 12 (sub-group 0, local id 12): ballot mask 0xffffffff, opportunistic group size 32
item 13 (sub-group 0, local id 13): ballot mask 0xffffffff, opportunistic group size 32
item 14 (sub-group 0, local id 14): ballot mask 0xffffffff, opportunistic group size 32
item 15 (sub-group 0, local id 15): ballot mask 0xffffffff, opportunistic group size 32
item 32 (sub-group 1, local id 0): ballot mask 0xffffffff, opportunistic group size 32
item 33 (sub-group 1, local id 1): ballot mask 0xffffffff, opportunistic group size 32
item 34 (sub-group 1, local id 2): ballot mask 0xffffffff, opportunistic group size 32
item 35 (sub-group 1, local id 3): ballot mask 0xffffffff, opportunistic group size 32
item 36 (sub-group 1, local id 4): ballot mask 0xffffffff, opportunistic group size 32
item 37 (sub-group 1, local id 5): ballot mask 0xffffffff, opportunistic group size 32
item 38 (sub-group 1, local id 6): ballot mask 0xffffffff, opportunistic group size 32
item 39 (sub-group 1, local id 7): ballot mask 0xffffffff, opportunistic group size 32
item 40 (sub-group 1, local id 8): ballot mask 0xffffffff, opportunistic group size 32
item 41 (sub-group 1, local id 9): ballot mask 0xffffffff, opportunistic group size 32
item 42 (sub-group 1, local id 10): ballot mask 0xffffffff, opportunistic group size 32
item 43 (sub-group 1, local id 11): ballot mask 0xffffffff, opportunistic group size 32
item 44 (sub-group 1, local id 12): ballot mask 0xffffffff, opportunistic group size 32
item 45 (sub-group 1, local id 13): ballot mask 0xffffffff, opportunistic group size 32
item 46 (sub-group 1, local id 14): ballot mask 0xffffffff, opportunistic group size 32
item 47 (sub-group 1, local id 15): ballot mask 0xffffffff, opportunistic group size 32
FAILED
sub-group size 16, converged first
expected: ballot mask 0xff, opportunistic group size 8
item 0 (sub-group 0, local id 0): ballot mask 0xffff, opportunistic group size 16
item 1 (sub-group 0, local id 1): ballot mask 0xffff, opportunistic group size 16
item 2 (sub-group 0, local id 2): ballot mask 0xffff, opportunistic group size 16
item 3 (sub-group 0, local id 3): ballot mask 0xffff, opportunistic group size 16
item 4 (sub-group 0, local id 4): ballot mask 0xffff, opportunistic group size 16
item 5 (sub-group 0, local id 5): ballot mask 0xffff, opportunistic group size 16
item 6 (sub-group 0, local id 6): ballot mask 0xffff, opportunistic group size 16
item 7 (sub-group 0, local id 7): ballot mask 0xffff, opportunistic group size 16
item 16 (sub-group 1, local id 0): ballot mask 0xffff, opportunistic group size 16
item 17 (sub-group 1, local id 1): ballot mask 0xffff, opportunistic group size 16
item 18 (sub-group 1, local id 2): ballot mask 0xffff, opportunistic group size 16
item 19 (sub-group 1, local id 3): ballot mask 0xffff, opportunistic group size 16
item 20 (sub-group 1, local id 4): ballot mask 0xffff, opportunistic group size 16
item 21 (sub-group 1, local id 5): ballot mask 0xffff, opportunistic group size 16
item 22 (sub-group 1, local id 6): ballot mask 0xffff, opportunistic group size 16
item 23 (sub-group 1, local id 7): ballot mask 0xffff, opportunistic group size 16
FAILED
sub-group size 32, converged first
expected: ballot mask 0xffff, opportunistic group size 16
item 0 (sub-group 0, local id 0): ballot mask 0xffffffff, opportunistic group size 32
item 1 (sub-group 0, local id 1): ballot mask 0xffffffff, opportunistic group size 32
item 2 (sub-group 0, local id 2): ballot mask 0xffffffff, opportunistic group size 32
item 3 (sub-group 0, local id 3): ballot mask 0xffffffff, opportunistic group size 32
item 4 (sub-group 0, local id 4): ballot mask 0xffffffff, opportunistic group size 32
item 5 (sub-group 0, local id 5): ballot mask 0xffffffff, opportunistic group size 32
item 6 (sub-group 0, local id 6): ballot mask 0xffffffff, opportunistic group size 32
item 7 (sub-group 0, local id 7): ballot mask 0xffffffff, opportunistic group size 32
item 8 (sub-group 0, local id 8): ballot mask 0xffffffff, opportunistic group size 32
item 9 (sub-group 0, local id 9): ballot mask 0xffffffff, opportunistic group size 32
item 10 (sub-group 0, local id 10): ballot mask 0xffffffff, opportunistic group size 32
item 11 (sub-group 0, local id 11): ballot mask 0xffffffff, opportunistic group size 32
item 12 (sub-group 0, local id 12): ballot mask 0xffffffff, opportunistic group size 32
item 13 (sub-group 0, local id 13): ballot mask 0xffffffff, opportunistic group size 32
item 14 (sub-group 0, local id 14): ballot mask 0xffffffff, opportunistic group size 32
item 15 (sub-group 0, local id 15): ballot mask 0xffffffff, opportunistic group size 32
item 32 (sub-group 1, local id 0): ballot mask 0xffffffff, opportunistic group size 32
item 33 (sub-group 1, local id 1): ballot mask 0xffffffff, opportunistic group size 32
item 34 (sub-group 1, local id 2): ballot mask 0xffffffff, opportunistic group size 32
item 35 (sub-group 1, local id 3): ballot mask 0xffffffff, opportunistic group size 32
item 36 (sub-group 1, local id 4): ballot mask 0xffffffff, opportunistic group size 32
item 37 (sub-group 1, local id 5): ballot mask 0xffffffff, opportunistic group size 32
item 38 (sub-group 1, local id 6): ballot mask 0xffffffff, opportunistic group size 32
item 39 (sub-group 1, local id 7): ballot mask 0xffffffff, opportunistic group size 32
item 40 (sub-group 1, local id 8): ballot mask 0xffffffff, opportunistic group size 32
item 41 (sub-group 1, local id 9): ballot mask 0xffffffff, opportunistic group size 32
item 42 (sub-group 1, local id 10): ballot mask 0xffffffff, opportunistic group size 32
item 43 (sub-group 1, local id 11): ballot mask 0xffffffff, opportunistic group size 32
item 44 (sub-group 1, local id 12): ballot mask 0xffffffff, opportunistic group size 32
item 45 (sub-group 1, local id 13): ballot mask 0xffffffff, opportunistic group size 32
item 46 (sub-group 1, local id 14): ballot mask 0xffffffff, opportunistic group size 32
item 47 (sub-group 1, local id 15): ballot mask 0xffffffff, opportunistic group size 32
FAILED
```

### Environment

- OS: Linux
- Target device and vendor: Intel(R) Data Center GPU Max 1550 (`intel_gpu_pvc`), Level Zero backend
- DPC++ version: Intel oneAPI 2026.1
```
$ icpx --version
Intel(R) oneAPI DPC++/C++ Compiler 2026.1.0 (2026.1.0.20260617)
Target: x86_64-unknown-linux-gnu
Thread model: posix
```
- Dependencies version: Intel compute-runtime 25.18.33578 (Level Zero driver 1.6.33578+77)

### Additional context

- The problem is present both when the work-items diverge before the first group operation, and when they all perform a group operation in converged control flow before diverging.
- The problem affects both sub-group sizes supported by the device (16 and 32).
- The problem is present both with ahead-of-time compilation (`-fsycl-targets=spir64_gen -Xs "-device pvc"`) and with just-in-time compilation (`-fsycl-targets=spir64`).

Reproduce the issue with the supplied test.cc on the Intel Data Center GPU Max 1550 at -O0 and -O2. Inspect sycl/include/sycl/ext/oneapi/experimental/fragment.hpp and the group_ballot implementation to trace the SPIR-V behavior. Done means returned work-items are excluded from both the ballot mask and opportunistic-group size at either optimization level.

Written by the indexing model from the issue text.