libigc.so.2: internal heap buffer overflow during SYCL kernel JIT corrupts adjacent glibc chunk metadata (DG2 / Arc A770)

## Environment

| Component | Version |
|---|---|
| Compiler | Intel oneAPI DPC++/C++ Compiler 2026.0.0 (2026.0.0.20260331) |
| IGC | `intel-igc-cm` 1.0.225.54093-1136~24.04 |
| OpenCL ICD | `intel-opencl-icd` 25.13.33276.22-1133~24.04 |
| Level Zero loader | `libze1` 1.21.9.0-1136~24.04 |
| Level Zero Intel GPU runtime | `libze-intel-gpu1` 25.13.33276.22-1133~24.04 |
| GPU | Intel Arc A770 Graphics (DG2, driver 1.6.33276+22) |
| OS | Ubuntu 24.04 LTS, kernel 6.17.0-19-generic |

## Summary

Under the normal glibc allocator, every SYCL `parallel_for` dispatch that
triggers IGC JIT compilation corrupts the heap by ~1–8 bytes past one of
IGC's internal buffers. The corruption is silent at the time of the write —
the overflow lands in the *bk* pointer of an adjacent free glibc chunk. The
process continues running normally until exit, when the destructor chain
walks the now-corrupt free list and the unlink invariant check fires:

```
free(): invalid pointer
```

or, depending on which free chunk's metadata was corrupted:

```
free(): corrupted unsorted chunks
```

…and the process SIGABRTs during shutdown.

## Why this looks "intermittent"

Whether the exit-time SIGABRT fires depends on heap layout — specifically,
whether IGC's overflow lands in a chunk that the exit-time destructor walks.
Any of the following can flip the bug on or off:

- Binary `.text` / `.bss` section size shifts (any code change)
- glibc main-arena tunable thresholds (e.g. `MALLOC_ARENA_MAX`, dynamic
`pthread_create` allocations)
- Process allocation history (thread pool init order, library load order)
- Tiny (sub-KiB) heap allocations elsewhere in the process

## Confirmation: two independent stack traces

Caught with `LD_PRELOAD=libtcmalloc.so.4` + GDB, where tcmalloc's
mmap-delegation makes the corruption manifest **during runtime** instead of
at exit. Both traces below show the corruption originating inside
`libigc.so.2`, not in the application. They were captured from two different
SYCL kernels with different allocation histories, to rule out an
application-side overflow.

### Trace 1 (kernel A: a simple element-wise kernel)

```
#0 __pthread_kill_implementation
#1 __pthread_kill_internal
#2 __GI_raise
#3 __GI_abort
#4 __libc_message
#5 malloc_printerr
#6 munmap_chunk
#7 __GI___libc_free (mem=0xeafabf80)
#8 ?? () from /opt/intel/oneapi/compiler/2026.0/lib/libigc.so.2
#9 IGC::IgcOclTranslationCtx::Translate
#10 ...SYCL kernel compile path...
```

### Trace 2 (kernel B: a different kernel, different alloc history)

```
#0 __pthread_kill_implementation
#1 __pthread_kill_internal
#2 __GI_raise
#3 __GI_abort
#4 __libc_message
#5 malloc_printerr ("double free or corruption (out)")
#6 _int_free_merge_chunk (av=0x7fffe..., p=0x..., size=3913525344)
#7 __GI___libc_free (mem=0xe9d732e0)
#8 ?? () from /opt/intel/oneapi/compiler/2026.0/lib/libigc.so.2
#9 IGC::IgcOclTranslationCtx::Translate
#10 ...
```

The `size=3913525344` (~3.6 GiB) value at frame 6 is clearly bogus — the
process heap is nowhere near that large. It is what has been written *into*
the adjacent free chunk's `mchunk_size` field by IGC's overflow.

## Why other tools didn't catch it

- **valgrind** is blocked because the Intel UR Level Zero adapter does driver
mmap operations on device memory regions that valgrind cannot instrument;
the process SIGSEGVs inside the L0 runtime before any application code runs.
- **icpx ASAN** is blocked because the Intel libomp/iomp5 thread stack isn't
recognized by ASAN's stack detector, producing immediate "wild pointer"
false positives at startup and garbled output.
- **glibc `MALLOC_CHECK_=3` + `MALLOC_PERTURB_=65`** runs cleanly because
glibc's check happens at *free* time on the chunk being freed; the
corruption is on a *different* chunk's metadata.
- **jemalloc `junk:true`** runs cleanly because jemalloc uses thread-local
arenas and size-class bins; IGC's overflow lands inside jemalloc's own
metadata, which jemalloc doesn't check on free.
- **tcmalloc** *does* catch it because tcmalloc routes large allocations
through `munmap_chunk` (glibc's invariant check), which fires on the
corrupted size value.
- **Electric Fence (`libefence`)** can't run because it `dlopen`s
page-guarded malloc against the entire process, which the L0 driver's
device buffer registration can't tolerate (page faults on driver-internal
pointers).

## Minimal reproducer

```cpp
#include

int main() {
sycl::queue q{sycl::gpu_selector_v};
int *p = sycl::malloc_device(256, q);

// Any parallel_for triggers IGC JIT on first dispatch
q.parallel_for(sycl::range<1>(256), [=](sycl::id<1> i) {
p[i] = (int)i.get(0);
}).wait();

sycl::free(p, q);
return 0; // SIGABRT here under glibc + specific layouts
}
```

Build:
```
icpx -fsycl -fsycl-targets=spir64 repro.cpp -o repro
```

Run under tcmalloc to surface the corruption immediately:
```
LD_PRELOAD=/usr/lib/x86_64-linux-gnu/libtcmalloc.so.4 ./repro
```

The reproducer is non-deterministic — repeated runs with different binary
layouts (e.g. adding an unused static char array of varying size) flip the
crash on and off, consistent with the heap-layout-sensitivity described
above.

## Suggested next steps for the IGC team

1. ASAN-build IGC itself (`libigc.so.2`) and run the reproducer under it.
The application binary does not need to be ASAN-built; only IGC.
2. Audit the kernel-translation buffer allocations in
`IGC::IgcOclTranslationCtx::Translate` (and the source-IR copy path it
calls) for `+1` / `+sizeof()` off-by-ones.
3. The bug is reproducible on DG2; please confirm whether the same issue
exists on Tiger Lake / Alder Lake iGPUs (which compile via the same IGC
code path).

## Application-side workaround (for other affected users)

We mitigate by (a) sending SIGKILL on process stop so the corrupt free-list
is never walked at teardown, (b) persisting the SYCL kernel cache so JIT is
skipped on subsequent starts, and (c) an opt-in ahead-of-time build
(`-fsycl-targets=spir64_gen -Xs "-device dg2"`) that eliminates IGC JIT
entirely. Happy to share exact config if useful.