diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml
index 98c7e488..07c37f37 100644
--- a/.github/workflows/ci.yml
+++ b/.github/workflows/ci.yml
@@ -15,9 +15,21 @@ permissions:
 
 jobs:
   self-runner:
-    continue-on-error: true
+    # Spelled out so the PR checks box says which driver stack this exercises: this job is
+    # the LTS stack (ONEAPI_LTS=1, self-hosted Aurora runner), while buildkite covers the
+    # rolling stack. Without an explicit name GitHub derives one from the matrix
+    # ("self-runner (ubuntu-latest, 1, x64)"), which says nothing about the stack and
+    # misreports the runner as ubuntu-latest — `os` is unused, `runs-on` is self-hosted.
+    name: LTS stack (self-hosted, Julia ${{ matrix.julia-version }})
     env:
       GITHUB_TOKEN: ${{ secrets.GITHUB_TOKEN }}
+      # Pin each parallel test worker to a distinct GPU tile instead of
+      # oversubscribing device 0 (see test/runtests.jl).
+      ONEAPI_TEST_SPREAD_GPUS: '1'
+      ONEAPI_LTS: '1'
+      # Synchronize after every command-list submission to work around the
+      # Aurora LTS NEO dropped-tail corruption (see lib/level-zero/cmdlist.jl).
+      ONEAPI_SYNC_EACH_SUBMISSION: '1'
     runs-on: [self-hosted, linux, X64]
     strategy:
       matrix:
@@ -36,7 +48,50 @@ jobs:
         with:
           version: ${{ matrix.julia-version }}
       - uses: julia-actions/cache@v3
+      # Build liboneapi_support.so from deps/src rather than using the registered
+      # oneAPI_Support_jll artifact, which lags behind the wrappers in this branch.
+      # Writes LocalPreferences.toml (plus a copy in test/, for Pkg.jl#2500) pointing
+      # oneAPI_Support_jll at the locally-built library. Not continue-on-error: if this
+      # fails we would silently fall back to the JLL and test the wrong library.
+      - name: Build the oneAPI support library
+        timeout-minutes: 90
+        run: julia --color=yes --project=deps deps/build_local.jl
       - uses: julia-actions/julia-buildpkg@latest
         continue-on-error: true
-      - uses: julia-actions/julia-runtest@latest
-        continue-on-error: true
+      # Must run after julia-buildpkg: Manifest.toml is gitignored, so the top-level
+      # project is not instantiated until then and `using oneAPI_Support_jll` would
+      # fail with "required but does not seem to be installed".
+      - name: Check the local support library is picked up
+        run: |
+          julia --color=yes --project=. -e '
+              using TOML, oneAPI_Support_jll
+              want = TOML.parsefile("LocalPreferences.toml")["oneAPI_Support_jll"]["liboneapi_support_path"]
+              got = oneAPI_Support_jll.liboneapi_support_path
+              @info "support library" want got
+              want == got || error("oneAPI_Support_jll is not using the locally-built library")'
+      # Disable AVX512-FP16 host codegen on the Aurora Sapphire Rapids nodes. Under concurrent
+      # oneMKL load the native AVX512-FP16 path silently miscomputes *host* Float16 (e.g. the
+      # GPUArrays `A .* B .+ c` broadcast reference), failing tests even though the GPU result
+      # is correct (single-process clean; MXCSR clean; only the native-FP16 path, not Float32).
+      # `-C native,-avx512fp16` routes Float16 through Float32 and propagates to the Pkg.test
+      # subprocess and its parallel workers via Base.julia_cmd(). `julia-runtest` cannot pass a
+      # cpu-target, so invoke Pkg.test() directly.
+      - name: Run tests (AVX512-FP16 disabled)
+        run: julia -C "native,-avx512fp16" --color=yes --project=. -e 'import Pkg; Pkg.test(coverage=true)'
+      # Report coverage from this job too, not just from buildkite. Every LTS-gated branch
+      # (the queue registry, the coalesced reduction path, the :khronos codegen selection,
+      # the drain-before-free) is unreachable on the rolling stack, so without this upload
+      # Codecov counts all of it as uncovered. `--code-coverage` propagates to the Pkg.test
+      # subprocess and its parallel workers through Base.julia_cmd(), as the cpu-target
+      # above does. Directories match the buildkite julia-coverage plugin's.
+      - uses: julia-actions/julia-processcoverage@v1
+        with:
+          directories: src,lib,examples
+      - uses: codecov/codecov-action@v5
+        with:
+          files: lcov.info
+          token: ${{ secrets.CODECOV_TOKEN }}
+          # Coverage reporting must not be able to fail the build: this job is the only
+          # signal for the LTS stack, and a Codecov outage or a missing token is not a
+          # test failure.
+          fail_ci_if_error: false
diff --git a/Project.toml b/Project.toml
index 20bf5ddc..98b423dd 100644
--- a/Project.toml
+++ b/Project.toml
@@ -1,7 +1,7 @@
 name = "oneAPI"
 uuid = "8f75cd03-7ff8-4ecb-9b8f-daf728133b1b"
 authors = ["Tim Besard <tim.besard@gmail.com>", "Alexis Montoison", "Michel Schanen <michel.schanen@gmail.com>"]
-version = "2.7.3"
+version = "2.8.0"
 
 [deps]
 AbstractFFTs = "621f4979-c628-5d54-868e-fcf4e3e8185c"
@@ -23,6 +23,7 @@ Printf = "de0858da-6303-5e67-8744-51eddeeeb8d7"
 Random = "9a3f8284-a2c9-5f02-9a11-845980a1fd5c"
 SPIRVIntrinsics = "71d1d633-e7e8-4a92-83a1-de8814b09ba8"
 SPIRV_LLVM_Backend_jll = "4376b9bf-cff8-51b6-bb48-39421dff0d0c"
+SPIRV_LLVM_Translator_jll = "4a5d46fc-d8cf-5151-a261-86b458210efb"
 SPIRV_Tools_jll = "6ac6d60f-d740-5983-97d7-a4482c0689f4"
 SparseArrays = "2f01184e-e22b-5df5-ae63-d93ebab69eaf"
 SpecialFunctions = "276daf66-3868-5448-9aa4-cd146d93841b"
@@ -47,6 +48,7 @@ PrecompileTools = "1"
 Preferences = "1"
 SPIRVIntrinsics = "1"
 SPIRV_LLVM_Backend_jll = "22"
+SPIRV_LLVM_Translator_jll = "21"
 SPIRV_Tools_jll = "2025.4.0"
 SpecialFunctions = "1.3, 2"
 StaticArrays = "1"
diff --git a/README.md b/README.md
index 778dd221..16be7180 100644
--- a/README.md
+++ b/README.md
@@ -73,6 +73,7 @@ Binary dependencies:
 - libigc: 2.34.4+1
 - gmmlib: 22.10.0+0
 - SPIRV_LLVM_Backend: 22.1.8+1
+- SPIRV_LLVM_Translator: 21.1.1+0
 - SPIRV_Tools: 2025.4.0+0
 - oneAPI_Support: 0.10.0+0 (oneMKL v2025.3.0)
 
@@ -154,7 +155,10 @@ julia> execute!(queue) do list
 Built on top of that, are kernel programming capabilities for executing Julia code on oneAPI
 accelerators. Device-side intrinsics are provided by
 [SPIRVIntrinsics.jl](https://github.com/JuliaGPU/SPIRVIntrinsics.jl), and code is compiled to
-SPIR-V using [LLVM's SPIR-V back-end](https://llvm.org/docs/SPIRVUsage.html):
+SPIR-V using [LLVM's SPIR-V back-end](https://llvm.org/docs/SPIRVUsage.html) — or, on Intel's
+LTS driver stack, the [Khronos SPIR-V
+translator](https://github.com/KhronosGroup/SPIRV-LLVM-Translator) (see [Intel's LTS driver
+stack](#intels-lts-driver-stack) below):
 
 ```julia
 julia> function kernel()
@@ -292,6 +296,38 @@ julia> oneArray([1.]) .+ 1
 │ error: Double type is not supported on this platform.
 ```
 
+### Intel's LTS driver stack
+
+Intel ships the Compute Runtime both as rolling releases and as a long-term-servicing (LTS)
+branch that large deployments — Aurora, for instance — stay on for years. oneAPI.jl targets
+the rolling stack by default. The LTS branch predates a number of driver and compiler fixes,
+some of which silently corrupt results rather than raise an error, so the package carries a
+set of workarounds behind an opt-in switch:
+
+```bash
+export ONEAPI_LTS=1
+```
+
+With it enabled, oneAPI.jl:
+
+- compiles kernels to SPIR-V with the [Khronos SPIR-V
+  translator](https://github.com/KhronosGroup/SPIRV-LLVM-Translator) instead of LLVM's
+  SPIR-V back-end, which the LTS NEO/IGC runtime does not accept;
+- disables `BFloat16`, which that SPIR-V stack cannot translate in generic kernels;
+- materializes strided inputs and uses a coalesced reduction kernel, working around an IGC
+  miscompile of non-coalesced reads that silently corrupts e.g. `sum(transpose(x))`;
+- drains command queues before freeing their buffers, since LTS NEO ignores the Level Zero
+  `BLOCKING_FREE` policy and would otherwise fault and ban the context.
+
+A separate switch, `ONEAPI_SYNC_EACH_SUBMISSION=1`, additionally synchronizes after every
+command-list submission. That works around a "dropped tail" corruption seen when several
+processes oversubscribe a single tile, at roughly a 3× throughput cost.
+
+Neither switch installs an LTS driver — they tell oneAPI.jl how to behave against one, and
+are typically combined with the local-toolchain configuration described below. Leave both
+unset on a rolling-release driver: the workarounds trade performance for correctness, and
+none of them are needed there. See the
+[documentation](https://juliagpu.github.io/oneAPI.jl/dev/lts/) for the full details.
 
 
 ## Development
diff --git a/docs/make.jl b/docs/make.jl
index fe6dd6f2..a5afba1d 100644
--- a/docs/make.jl
+++ b/docs/make.jl
@@ -27,6 +27,7 @@ makedocs(
             "Memory Management" => "memory.md",
             "Device Intrinsics" => "device.md",
             "Performance Guide" => "usage/performance.md",
+            "Intel LTS Driver Stack" => "lts.md",
         ],
         "API Reference" => [
             "Overview" => "api.md",
diff --git a/docs/src/index.md b/docs/src/index.md
index ad959992..c8b62df9 100644
--- a/docs/src/index.md
+++ b/docs/src/index.md
@@ -17,7 +17,8 @@ Direct use of `oneAPI`-specific macros (like `@oneapi`) and types (like `oneArra
 
 - **High-level Array Abstractions**: `oneArray` type fully implementing the `GPUArrays.jl` interface.
 - **Kernel Programming**: Execute custom kernels written in Julia on Intel GPUs, compiled to
-  SPIR-V through LLVM's SPIR-V back-end.
+  SPIR-V through LLVM's SPIR-V back-end (or the Khronos SPIR-V translator on Intel's
+  [LTS driver stack](lts.md)).
 - **KernelAbstractions.jl Backend**: The `oneAPIBackend` back-end runs portable kernels on
   Intel GPUs unmodified.
 - **Level Zero Integration**: Low-level access to the Level Zero API via the `oneL0` submodule.
diff --git a/docs/src/installation.md b/docs/src/installation.md
index 59cf36ae..4f60a13a 100644
--- a/docs/src/installation.md
+++ b/docs/src/installation.md
@@ -45,6 +45,21 @@ julia> devices()
 julia> device!(1) # Select the first available device
 ```
 
+## Intel LTS Driver Stack
+
+Large deployments (for example Aurora) run Intel's long-term-servicing branch of the
+Compute Runtime rather than a rolling release. That branch needs a number of workarounds,
+which oneAPI.jl provides behind an opt-in switch:
+
+```bash
+export ONEAPI_LTS=1
+```
+
+Among other things this compiles kernels with the Khronos SPIR-V translator instead of
+LLVM's SPIR-V back-end. See [Intel LTS Driver Stack](lts.md) for what else changes and what
+it costs. The switch does not install an LTS driver — combine it with the system-library
+configuration below.
+
 ## Using System Libraries (Advanced)
 
 !!! warning
diff --git a/docs/src/lts.md b/docs/src/lts.md
new file mode 100644
index 00000000..6251b10c
--- /dev/null
+++ b/docs/src/lts.md
@@ -0,0 +1,155 @@
+# Intel LTS driver stack
+
+Intel ships the Compute Runtime in two lines: frequent *rolling* releases, and a
+long-term-servicing (LTS) branch that large deployments stay on for years — for example
+the NEO/IGC LTS stack on [Aurora](https://www.alcf.anl.gov/aurora). oneAPI.jl targets the
+rolling stack by default.
+
+The LTS branch predates a number of driver and compiler fixes that landed in rolling. Left
+alone, some of those defects are not merely inconvenient: they silently corrupt results, or
+get the Level Zero context banned so that every later submission fails. oneAPI.jl therefore
+carries a set of workarounds, kept behind a single opt-in switch so that the default
+(rolling) path is completely unaffected.
+
+## Enabling
+
+Set the environment variable before loading oneAPI.jl:
+
+```bash
+export ONEAPI_LTS=1
+```
+
+`1`, `true`, `yes` and `on` enable it; `0`, `false`, `no` and `off` disable it. An
+unrecognized value warns and falls back to the default (disabled), rather than silently
+reading as off.
+
+You can check which mode is active from Julia:
+
+```julia
+julia> using oneAPI
+
+julia> oneAPI.oneL0.LTS[]
+true
+```
+
+!!! note
+    This flag does not install or select an LTS driver. oneAPI.jl ships pinned rolling
+    NEO artifacts; on an LTS system you point the package at the system libraries instead
+    (see the "Using System Libraries" section of [Installation](installation.md)).
+    `ONEAPI_LTS=1` tells oneAPI.jl to *compile and behave* for that driver.
+
+## What changes
+
+### SPIR-V code generation
+
+Julia kernels are compiled to SPIR-V. By default oneAPI.jl uses [LLVM's SPIR-V
+back-end](https://llvm.org/docs/SPIRVUsage.html) (`SPIRV_LLVM_Backend_jll`). The LTS
+NEO/IGC runtime does not accept that output, so with `ONEAPI_LTS=1` the package switches to
+the [Khronos SPIR-V translator](https://github.com/KhronosGroup/SPIRV-LLVM-Translator)
+(`SPIRV_LLVM_Translator_jll`).
+
+Both tools are dependencies of oneAPI.jl and are resolved lazily by GPUCompiler, so the
+choice costs nothing on either path and needs no reconfiguration of your environment.
+`oneAPI.versioninfo()` reports both, and `@device_code_spirv` will show which one produced a
+given module in its `Generator:` line.
+
+### BFloat16 is unavailable
+
+The LTS SPIR-V stack cannot translate the LLVM `bfloat` type in generic kernels: a kernel
+that merely keeps a `BFloat16` value (`clamp!`, say) fails with an `InvalidIRError`, and
+declaring the `SPV_KHR_bfloat16` extension crashes the LTS runtime outright. On LTS the
+compiler is therefore configured with `supports_bfloat16 = false`, regardless of what the
+device reports — `oneAPI._device_supports_bfloat16()` is a *hardware* capability check and
+does not capture this limitation.
+
+Consequently `BFloat16` is dropped from the element types exercised by the test suite, and
+`examples/bfloat16.jl` exits early with a message. Other floating-point types are unaffected.
+
+### Reductions over strided inputs
+
+The LTS IGC miscompiles non-coalesced (strided) global reads inside the reduction kernel,
+silently producing wrong results — no error, just bad numbers. Reading an array along a
+non-contiguous axis is enough to trigger it, so `sum(transpose(x))`, `a == transpose(b)` and
+`ishermitian(x)` are all affected. Elementwise copies are not.
+
+Two mitigations apply on LTS:
+
+- An input that is not densely laid out (a transposed, permuted or otherwise strided view,
+  or a broadcast containing one) is materialized into a dense `oneArray` before the
+  reduction runs. This costs an extra allocation and copy.
+- A reduction that keeps the contiguous leading dimension (`sum(A; dims=2)` and friends) is
+  routed to a coalesced kernel that assigns one work-item per output slice, so neighbouring
+  lanes read neighbouring memory. Reductions with few output slices get less parallelism
+  than the default workgroup-per-slice kernel, but stay correct.
+
+Reductions that include dimension 1 (`dims=(1,3)`, or a full reduction) keep a contiguous
+innermost axis and use the normal kernel.
+
+### Frees are synchronized against in-flight work
+
+NEO LTS advertises `ZE_extension_memory_free_policies` but does not honor its
+`BLOCKING_FREE` policy: it unmaps an allocation immediately, even with work in flight that
+references it. A garbage-collected free of a dead array whose last kernel has not retired
+then faults on the GPU, which bans the kernel context — after which *every* later submission
+fails with `ZE_RESULT_ERROR_UNKNOWN`.
+
+On LTS, oneAPI.jl keeps a registry of the command queues in use and drains those that could
+reference a buffer before freeing it. Command queues are likewise drained before being
+destroyed; a queue still busy after 10 s is deliberately leaked, since destroying it would
+trigger the very fault the drain prevents. The visible cost is that a GC-driven free can
+block until outstanding work completes.
+
+### Optional: synchronize after every submission
+
+Under heavy multi-process oversubscription of a single tile, a whole-queue
+`zeCommandQueueSynchronize` on the LTS stack does not reliably retire the tail of an
+earlier, separately submitted command list. The result is a silent *dropped tail*: the last
+work-items of a kernel, or the last elements of a copy, never land.
+
+Synchronizing after every submission eliminates it, at roughly a 3× throughput cost. It is
+off by default and enabled with:
+
+```bash
+export ONEAPI_SYNC_EACH_SUBMISSION=1
+```
+
+It can also be controlled at runtime:
+
+```julia
+oneL0.sync_each_submission()            # query
+oneL0.sync_each_submission!(true)       # set, returns the previous value
+
+oneL0.sync_each_submission(false) do    # scoped, restores afterwards
+    # ...
+end
+```
+
+!!! warning
+    A *submit-then-signal* pattern — work submitted with a wait event that is only signaled
+    after submission returns — deadlocks with this enabled, because `execute!` blocks in the
+    synchronize before the gating event can be signaled. No high-level oneAPI.jl code path
+    submits event-gated work, but hand-written Level Zero code may; wrap those regions in
+    `oneL0.sync_each_submission(false) do ... end`.
+
+This workaround is independent of `ONEAPI_LTS` and can be enabled on its own.
+
+## Caveats
+
+- **Precompilation.** The precompilation workload warms up whichever SPIR-V tool
+  `ONEAPI_LTS` selects *at precompile time*, and Julia does not invalidate the cache when an
+  environment variable changes. After flipping the variable, run `Pkg.precompile()` to
+  refresh it. This only affects first-call latency, never correctness.
+- **No effect on the rolling stack.** Every workaround above is gated on the switch, so with
+  `ONEAPI_LTS` unset oneAPI.jl behaves exactly as it does upstream. The workarounds trade
+  performance for correctness and should not be enabled on a driver that does not need them.
+
+## Running the test suite
+
+The test suite reads the same switch and skips the cases the LTS stack cannot support (the
+`BFloat16` element type, the BFloat16 example). On a multi-tile node, pinning each test
+worker to its own tile avoids the oversubscription that provokes the dropped-tail
+corruption:
+
+```bash
+ONEAPI_LTS=1 ONEAPI_TEST_SPREAD_GPUS=1 julia --project=. -e 'import Pkg; Pkg.test()'
+```
diff --git a/docs/src/troubleshooting.md b/docs/src/troubleshooting.md
index 936b435c..7ff19357 100644
--- a/docs/src/troubleshooting.md
+++ b/docs/src/troubleshooting.md
@@ -36,6 +36,36 @@ Some Intel GPUs (especially integrated graphics) lack native hardware support fo
 - When several processes share one GPU, give each a budget with `ONEAPI_MEMORY_LIMIT`; see
   [Memory Management](memory.md).
 
+### Silently wrong results, or every submission failing
+
+**Symptom**: on a system running Intel's long-term-servicing (LTS) Compute Runtime — such
+as Aurora — reductions over transposed or otherwise strided arrays return wrong values
+without raising an error, kernels drop their last work-items, or every submission starts
+failing with `ZE_RESULT_ERROR_UNKNOWN` after a garbage collection.
+
+**Solution**:
+Enable the LTS workarounds and restart Julia:
+
+```bash
+export ONEAPI_LTS=1
+```
+
+See [Intel LTS Driver Stack](lts.md) for the individual defects, the additional
+`ONEAPI_SYNC_EACH_SUBMISSION=1` switch for oversubscribed tiles, and the performance
+trade-offs involved. Do not enable these on a rolling-release driver, which does not need
+them.
+
+### "InvalidIRError" for a BFloat16 kernel
+
+**Symptom**: a kernel using `BFloat16` fails to compile even though the device reports
+BFloat16 support.
+
+**Solution**:
+The LTS SPIR-V stack cannot translate the LLVM `bfloat` type in generic kernels; the
+device-level check (`oneAPI._device_supports_bfloat16()`) reports hardware capability and
+does not capture this. Use `Float16` or `Float32` on that stack — see
+[Intel LTS Driver Stack](lts.md).
+
 ## Debugging
 
 ### Validation Layer
diff --git a/examples/bfloat16.jl b/examples/bfloat16.jl
index 3fd88da9..581bd0f3 100644
--- a/examples/bfloat16.jl
+++ b/examples/bfloat16.jl
@@ -14,6 +14,15 @@ if !bfloat16_supported
     exit()
 end
 
+# _device_supports_bfloat16() reflects hardware conversion support, but the Aurora LTS
+# SPIR-V stack (Khronos translator + NEO/IGC) cannot translate the `bfloat` LLVM type at
+# all: even a load/store forces SPV_KHR_bfloat16, which the LTS runtime rejects. So bf16
+# kernels can't be compiled there regardless of device support.
+if oneAPI.oneL0.LTS[]
+    @info "BFloat16 kernels are unsupported on the Aurora LTS SPIR-V stack, skipping."
+    exit()
+end
+
 # Conversions: Core.BFloat16 in Julia 1.12 may not have Float32 constructors yet
 float32_to_bf16(x::Float32) = reinterpret(Core.BFloat16, (reinterpret(UInt32, x) >> 16) % UInt16)
 bf16_to_float32(x::Core.BFloat16) = reinterpret(Float32, UInt32(reinterpret(UInt16, x)) << 16)
diff --git a/lib/level-zero/cmdlist.jl b/lib/level-zero/cmdlist.jl
index 133d2551..e464c639 100644
--- a/lib/level-zero/cmdlist.jl
+++ b/lib/level-zero/cmdlist.jl
@@ -47,8 +47,58 @@ function ZeCommandList(f::Base.Callable, args...; kwargs...)
     return list
 end
 
-execute!(queue::ZeCommandQueue, lists::Vector{ZeCommandList}, fence=nothing) =
-    zeCommandQueueExecuteCommandLists(queue, length(lists), lists, something(fence, C_NULL))
+# Opt-in workaround for the Aurora LTS NEO stack (set ONEAPI_SYNC_EACH_SUBMISSION=1).
+# Under heavy multi-process oversubscription of a single tile, a whole-queue
+# `zeCommandQueueSynchronize` does not reliably retire the tail of an earlier,
+# separately-submitted command list — producing silent "dropped tail" corruption (the
+# last work-item of a kernel / last element of a copy is missing). See
+# docs/src/lts.md. Synchronizing after *every* submission eliminates it, at a large
+# throughput cost (~3x), so it is off by default and only enabled when correctness under
+# oversubscription matters more than speed.
+const SYNC_EACH_SUBMISSION = Ref{Bool}(false)
+
+"""
+    sync_each_submission() -> Bool
+
+Whether [`execute!`](@ref) follows every command-list submission with a full
+`zeCommandQueueSynchronize` (the Aurora LTS "dropped tail" workaround). See
+[`sync_each_submission!`](@ref).
+"""
+sync_each_submission() = SYNC_EACH_SUBMISSION[]
+
+"""
+    sync_each_submission!(enable::Bool) -> Bool
+
+Enable or disable synchronizing after every submission, returning the previous setting.
+Initialized from the `ONEAPI_SYNC_EACH_SUBMISSION` environment variable.
+"""
+function sync_each_submission!(enable::Bool)
+    old = SYNC_EACH_SUBMISSION[]
+    SYNC_EACH_SUBMISSION[] = enable
+    return old
+end
+
+"""
+    sync_each_submission(f, enable::Bool)
+
+Run `f()` with the workaround temporarily set to `enable`, restoring the previous setting
+afterwards. Use this for submit-then-signal patterns that would otherwise deadlock, where a
+synchronize is forced before the event gating the submitted work is signaled.
+"""
+function sync_each_submission(f::Base.Callable, enable::Bool)
+    old = sync_each_submission!(enable)
+    return try
+        f()
+    finally
+        sync_each_submission!(old)
+    end
+end
+
+function execute!(queue::ZeCommandQueue, lists::Vector{ZeCommandList}, fence = nothing)
+    r = zeCommandQueueExecuteCommandLists(queue, length(lists), lists, something(fence, C_NULL))
+    sync_each_submission() && synchronize(queue)
+    return r
+end
 
 """
     execute!(queue::ZeCommandQueue, ...) do list
diff --git a/lib/level-zero/cmdqueue.jl b/lib/level-zero/cmdqueue.jl
index 3004a145..75ddb63b 100644
--- a/lib/level-zero/cmdqueue.jl
+++ b/lib/level-zero/cmdqueue.jl
@@ -2,6 +2,12 @@
 
 export ZeCommandQueue, synchronize
 
+# Bound on how long the LTS drain-before-destroy finalizer waits for in-flight work before
+# giving up and leaking the queue (see the finalizer below). 10 s comfortably covers any
+# kernel still legitimately running at finalization while keeping a never-signaled event
+# from hanging GC or process exit forever.
+const FINALIZER_SYNC_TIMEOUT_NS = UInt64(10_000_000_000)
+
 mutable struct ZeCommandQueue
     handle::ze_command_queue_handle_t
 
@@ -20,7 +26,30 @@ mutable struct ZeCommandQueue
         zeCommandQueueCreate(ctx, dev, desc_ref, handle_ref)
         obj = new(handle_ref[], ctx, dev, ordinal)
         finalizer(obj) do obj
+            if LTS[]
+                # the queue may still have work in flight (nothing requires a task to
+                # synchronize before dying), and zeCommandQueueDestroy does not wait for
+                # it: on the LTS NEO stack the still-running work then faults as soon as
+                # a referenced allocation is freed, getting the context banned. drain the
+                # queue first; unchecked, as sync on a banned context returns an error.
+                #
+                # Bounded wait, not typemax(UInt64): event-gated work whose event is never
+                # signaled (a task that submits work with a wait event and dies before
+                # signaling it) would make an infinite wait hang the finalizer forever, and
+                # with it GC and process exit. On timeout, leak the queue deliberately —
+                # destroying it now would trigger the very fault+ban this drain prevents.
+                if unchecked_zeCommandQueueSynchronize(obj, FINALIZER_SYNC_TIMEOUT_NS) == RESULT_NOT_READY
+                    @warn "Leaking a command queue still busy after $(FINALIZER_SYNC_TIMEOUT_NS ÷ 1_000_000_000)s to avoid blocking finalization (event-gated work whose event was never signaled?)" maxlog = 1
+                    return
+                end
+            end
             zeCommandQueueDestroy(obj)
+            if LTS[]
+                # mark the queue as destroyed: it can still be weakly reachable (e.g. from
+                # the queue registry used by `synchronize_all_queues`), and synchronizing a
+                # destroyed handle crashes in the driver.
+                obj.handle = ze_command_queue_handle_t(C_NULL)
+            end
         end
         obj
     end
diff --git a/lib/level-zero/oneL0.jl b/lib/level-zero/oneL0.jl
index 900f80de..948a213a 100644
--- a/lib/level-zero/oneL0.jl
+++ b/lib/level-zero/oneL0.jl
@@ -137,10 +137,41 @@ const functional = Ref{Bool}(false)
 const validation_layer = Ref{Bool}()
 const parameter_validation = Ref{Bool}()
 
+# Master switch for the Intel LTS-stack workarounds (driver/IGC quirks on the Aurora
+# LTS NEO 25.18 stack): the SPIR-V translator codegen path, strided-reduction
+# materialization, and command-queue drain-before-free. All such code paths are gated
+# on `LTS[]`, so with it disabled the package behaves like the upstream rolling stack.
+# Default OFF (rolling stack), matching upstream: an LTS deployment opts in with
+# `ONEAPI_LTS=1`. On Aurora that variable is set in the environment, and the GitHub
+# Actions self-hosted (Aurora) runner sets it in ci.yml; the rolling-stack buildkite
+# runner gets the default and thus actually exercises the non-LTS (:llvm back-end) path.
+const LTS = Ref{Bool}(false)
+
+# Parse a boolean-valued environment variable, accepting the same spellings for every
+# oneAPI flag. Returns `default` when the variable is unset or empty. Warns (and returns
+# `default`) on an unrecognized value, so a deployment that writes `ONEAPI_SYNC_EACH_SUBMISSION=on`
+# — valid for the sibling `ONEAPI_LTS` — no longer silently reads as off, and a typo like
+# `ONEAPI_LTS=treu` no longer silently disables the LTS path.
+function parse_env_bool(name::AbstractString, default::Bool)
+    haskey(ENV, name) || return default
+    val = lowercase(strip(ENV[name]))
+    isempty(val) && return default
+    val in ("1", "true", "yes", "on") && return true
+    val in ("0", "false", "no", "off") && return false
+    @warn "Ignoring unrecognized boolean value for $name; using default" value = ENV[name] default
+    return default
+end
+
 function __init__()
     precompiling = ccall(:jl_generating_output, Cint, ()) != 0
     precompiling && return
 
+    # Resolve the LTS master switch up front, before the driver-availability early
+    # returns below: it gates codegen and behavior and must be set even on hosts
+    # without a functional GPU. Default off (rolling stack); an LTS deployment such as
+    # Aurora opts in with ONEAPI_LTS=1.
+    LTS[] = parse_env_bool("ONEAPI_LTS", false)
+
     if Sys.iswindows()
         if Libdl.dlopen(libze_loader; throw_error=false) === nothing
             @error "The oneAPI Level Zero loader was not found. Please ensure the Intel GPU drivers are installed."
@@ -199,6 +230,7 @@ function __init__()
 
     validation_layer[] = parse(Bool, get(ENV, "ZE_ENABLE_VALIDATION_LAYER", "false"))
     parameter_validation[] = parse(Bool, get(ENV, "ZE_ENABLE_PARAMETER_VALIDATION", "false"))
+    return sync_each_submission!(parse_env_bool("ONEAPI_SYNC_EACH_SUBMISSION", false))
 end
 
 end
diff --git a/src/compiler/compilation.jl b/src/compiler/compilation.jl
index 928c1de9..b49d17b2 100644
--- a/src/compiler/compilation.jl
+++ b/src/compiler/compilation.jl
@@ -75,11 +75,14 @@ function GPUCompiler.finish_ir!(job::oneAPICompilerJob, mod::LLVM.Module,
     flatten_nested_insertvalue!(mod)
 
     # When the device supports BFloat16 but the SPIR-V runtime doesn't accept
-    # SPV_KHR_bfloat16, lower all bfloat types to i16 so the translator can
+    # SPV_KHR_bfloat16, lower all bfloat types to i16 so the back-end/translator can
     # handle the module without the extension. Both conditions are read from the
     # (device-independent) compiler config so this stays valid without a live
     # device: `_compiler_config` sets `supports_bfloat16` from the device and adds
     # the `SPV_KHR_bfloat16` extension iff the driver's SPIR-V runtime accepts it.
+    # Rolling stack only: bf16 is forced off on the LTS stack
+    # (`supports_bfloat16 = false`, see `_compiler_config`), so this pass never
+    # runs there.
     target = job.config.target
     if @static(isdefined(Core, :BFloat16) && isdefined(LLVM, :BFloatType)) &&
             target.supports_bfloat16 && !occursin("SPV_KHR_bfloat16", target.extensions)
@@ -324,26 +327,39 @@ end
     properties = oneL0.module_properties(dev)
     supports_fp16 = properties.fp16flags & oneL0.ZE_DEVICE_MODULE_FLAG_FP16 == oneL0.ZE_DEVICE_MODULE_FLAG_FP16
     supports_fp64 = properties.fp64flags & oneL0.ZE_DEVICE_MODULE_FLAG_FP64 == oneL0.ZE_DEVICE_MODULE_FLAG_FP64
-    # Allow BFloat16 in IR if the device supports it (even if the SPIR-V runtime doesn't
-    # advertise the extension). We lower bfloat→i16 in finish_ir! when needed.
-    supports_bfloat16 = _device_supports_bfloat16(dev)
-
-    # SPIR-V extensions the LLVM back-end may emit. Declaring them permits the
-    # corresponding instructions during translation: without
-    # SPV_EXT_shader_atomic_float_add, floating-point atomic operations fail to
-    # translate ("The atomic float instruction requires ... SPV_EXT_shader_atomic_float_add").
+
+    # SPIR-V codegen path. The Aurora LTS NEO/IGC runtime only accepts SPIR-V from the
+    # Khronos translator; the rolling stack uses the LLVM SPIR-V back-end. GPUCompiler picks
+    # the tool from the target's `backend` field and loads the JLL lazily, so both can be
+    # listed as deps and the choice is made here at compile time. Either way the extensions
+    # below have to be declared explicitly: without SPV_EXT_shader_atomic_float_add,
+    # floating-point atomic operations fail to translate ("The atomic float instruction
+    # requires ... SPV_EXT_shader_atomic_float_add").
+    # TODO: emit printf format strings in constant memory
     extensions = String[
         "SPV_EXT_relaxed_printf_string_address_space",
         "SPV_EXT_shader_atomic_float_add",
     ]
-    # Only add the SPIR-V extension if the runtime actually supports it
-    if _driver_supports_bfloat16_spirv(dev)
-        push!(extensions, "SPV_KHR_bfloat16")
+    if oneL0.LTS[]
+        backend = :khronos
+        # The LTS SPIR-V stack cannot codegen native bfloat in generic kernels (clamp! ->
+        # InvalidIRError, and SPV_KHR_bfloat16 segfaults NEO), so keep bf16 off entirely on
+        # this path; the test suite gates bf16 eltypes on `!oneL0.LTS[]` to match.
+        supports_bfloat16 = false
+    else
+        backend = :llvm
+        # Allow BFloat16 in IR if the device supports it (even if the SPIR-V runtime doesn't
+        # advertise the extension); finish_ir! lowers bfloat->i16 when the extension is absent.
+        supports_bfloat16 = _device_supports_bfloat16(dev)
+        # Only declare the SPIR-V extension when the runtime actually accepts it.
+        if _driver_supports_bfloat16_spirv(dev)
+            push!(extensions, "SPV_KHR_bfloat16")
+        end
     end
     extensions_str = join(map(ext -> "+$ext", extensions), ",")
 
     # create GPUCompiler objects
-    target = SPIRVCompilerTarget(; extensions=extensions_str, supports_fp16, supports_fp64, supports_bfloat16, kwargs...)
+    target = SPIRVCompilerTarget(; backend, extensions = extensions_str, supports_fp16, supports_fp64, supports_bfloat16, kwargs...)
     params = oneAPICompilerParams()
     CompilerConfig(target, params; kernel, name, always_inline)
 end
diff --git a/src/compiler/precompile.jl b/src/compiler/precompile.jl
index 4a8493c6..3d5567d8 100644
--- a/src/compiler/precompile.jl
+++ b/src/compiler/precompile.jl
@@ -3,9 +3,16 @@ using PrecompileTools: @compile_workload
 # Warm up the GPUCompiler -> SPIR-V pipeline during precompilation so the first real
 # `zefunction` call is cheap. This doesn't need a GPU: SPIR-V codegen is device-independent
 # (the target only carries version/extension/capability knobs), and the workload never
-# launches anything. It is gated on the SPIR-V LLVM back-end being available so the package
-# still precompiles on platforms/toolchains without it.
-if SPIRV_LLVM_Backend_jll.is_available()
+# launches anything. It is gated on the SPIR-V tool for the selected codegen path being
+# available so the package still precompiles on platforms/toolchains without it.
+#
+# The workload mirrors the runtime backend choice (see `_compiler_config`): the LTS stack
+# compiles through the Khronos translator, the rolling stack through the LLVM SPIR-V
+# back-end. `oneL0.__init__` has not run at precompile time, so resolve ONEAPI_LTS
+# directly from the environment with the same parser.
+const _precompile_lts = oneL0.parse_env_bool("ONEAPI_LTS", false)
+if _precompile_lts ? SPIRV_LLVM_Translator_jll.is_available() :
+        SPIRV_LLVM_Backend_jll.is_available()
     @compile_workload begin
         let
             function _precompile_kernel(a)
@@ -16,8 +23,11 @@ if SPIRV_LLVM_Backend_jll.is_available()
             # Build a device-independent compiler config. `_compiler_config` normally derives
             # these knobs from the device; here we use conservative, portable defaults (the
             # workload only exercises the pipeline, it does not target a specific device).
-            target = SPIRVCompilerTarget(; extensions="", supports_fp16=true,
-                                          supports_fp64=true, supports_bfloat16=false)
+            target = SPIRVCompilerTarget(;
+                backend = (_precompile_lts ? :khronos : :llvm),
+                extensions = "", supports_fp16 = true,
+                supports_fp64 = true, supports_bfloat16 = false,
+            )
             params = oneAPICompilerParams()
             config = CompilerConfig(target, params; kernel=true, name=nothing,
                                     always_inline=false)
diff --git a/src/context.jl b/src/context.jl
index 88035e74..f625a40d 100644
--- a/src/context.jl
+++ b/src/context.jl
@@ -224,10 +224,103 @@ function global_queue(ctx::ZeContext, dev::ZeDevice)
     # NOTE: dev purposefully does not default to context() or device() to stress that
     #       objects should track ownership, and not rely on implicit global state.
     get!(task_local_storage(), (:ZeCommandQueue, ctx, dev)) do
-        ZeCommandQueue(ctx, dev; flags = oneL0.ZE_COMMAND_QUEUE_FLAG_IN_ORDER)
+        queue = ZeCommandQueue(ctx, dev; flags = oneL0.ZE_COMMAND_QUEUE_FLAG_IN_ORDER)
+        register_queue!(ctx, dev, queue)
     end
 end
 
+# Register `queue` as a queue targeting (ctx, dev) so `synchronize_all_queues`/`release`
+# can find and drain it before freeing buffers whose in-flight work it may still reference.
+# EVERY queue that becomes a task's active queue must go through here — not just the one
+# `global_queue` creates but also the replacement `KA.priority!` installs — or the
+# unregistered queue's in-flight work can outlive a freed buffer (a use-after-free that
+# faults and bans the context on the LTS NEO stack). Only the LTS stack maintains the
+# registry; on the rolling stack this is a no-op. Returns `queue`.
+function register_queue!(ctx::ZeContext, dev::ZeDevice, queue::ZeCommandQueue)
+    oneL0.LTS[] || return queue
+    # disable finalizers while mutating the registry: a GC-driven finalizer on this
+    # task could call back into `synchronize_all_queues` (the lock is reentrant) and
+    # observe/mutate the registry mid-update.
+    GC.enable_finalizers(false)
+    try
+        @lock queue_registry_lock begin
+            push!(
+                get!(Vector{Tuple{WeakRef, ZeCommandQueue}}, queue_registry, (ctx, dev)),
+                (WeakRef(current_task()), queue)
+            )
+        end
+    finally
+        GC.enable_finalizers(true)
+    end
+    return queue
+end
+
+# Registry of all queues created through `global_queue`, across tasks. Buffers can be
+# freed from any task (GC finalizers), so `release` needs to be able to find the queues
+# that may still have work in flight referencing the buffer; queues themselves are
+# cached task-locally and would otherwise be unreachable from the finalizing task.
+#
+# Entries reference the queue *strongly*: the GC clears WeakRefs to a dead queue in the
+# same cycle that queues its finalizer, i.e., before the finalizer runs, so a WeakRef
+# would hide the queue from `release` exactly when its in-flight work still references
+# buffers about to be freed. The owning task is tracked weakly instead: queues are
+# task-local, so once their task is dead no new work can reach them, and the entry can
+# be dropped (allowing the queue to be finalized) after a final synchronize.
+const queue_registry_lock = ReentrantLock()
+const queue_registry = Dict{Tuple{ZeContext, ZeDevice}, Vector{Tuple{WeakRef, ZeCommandQueue}}}()
+
+# synchronize all known queues that target the given context (and device, if specified),
+# i.e., all queues whose in-flight work could possibly reference an allocation that is
+# about to be freed.
+function synchronize_all_queues(ctx::ZeContext, dev::Union{ZeDevice, Nothing})
+    # only the LTS stack populates the queue registry (see `global_queue`); on the
+    # rolling stack this is a no-op and `release` frees directly.
+    oneL0.LTS[] || return
+    queues = ZeCommandQueue[]
+    stale = Tuple{WeakRef, ZeCommandQueue}[]
+    GC.enable_finalizers(false)
+    try
+        @lock queue_registry_lock begin
+            for ((qctx, qdev), entries) in queue_registry
+                qctx == ctx || continue
+                (dev === nothing || qdev == dev) || continue
+                for entry in entries
+                    (task, queue) = entry
+                    queue.handle == C_NULL && continue  # finalized, handle destroyed
+                    push!(queues, queue)
+                    # entries whose task was already dead at this point cannot
+                    # receive new work, so they are safe to retire after the sync
+                    if task.value === nothing || istaskdone(task.value::Task)
+                        push!(stale, entry)
+                    end
+                end
+            end
+        end
+        # synchronize outside the lock: this can block for as long as a kernel runs,
+        # and finalizers running concurrently also need to take the lock. Keep
+        # finalizers disabled so none of the collected queues can be destroyed
+        # between collection and synchronization.
+        for queue in queues
+            oneL0.synchronize(queue)
+        end
+        # retire drained queues of dead tasks, allowing them to be finalized (the
+        # finalizer synchronizes once more before destroying the queue, in case
+        # the queue is dropped through other means).
+        if !isempty(stale)
+            @lock queue_registry_lock begin
+                for ((qctx, qdev), entries) in queue_registry
+                    qctx == ctx || continue
+                    (dev === nothing || qdev == dev) || continue
+                    filter!(entry -> !any(s -> s === entry, stale), entries)
+                end
+            end
+        end
+    finally
+        GC.enable_finalizers(true)
+    end
+    return
+end
+
 """
     synchronize()
 
diff --git a/src/mapreduce.jl b/src/mapreduce.jl
index 6813fdca..7fa8a74d 100644
--- a/src/mapreduce.jl
+++ b/src/mapreduce.jl
@@ -109,14 +109,68 @@ function partial_mapreduce_device(f, op, neutral, maxitems, Rreduce, Rother, R,
     return
 end
 
+# Coalesced reduction for when the contiguous leading dimension is NOT reduced (the reduced
+# axes are strided). One work-item per output slice (Rother element), grid-strided; each
+# serially reduces over Rreduce. Consecutive work-items map to consecutive output slices,
+# which are consecutive in memory, so global reads are coalesced across lanes — the access
+# pattern a `dims=1` reduction already uses. On the Aurora LTS stack the workgroup-per-slice
+# kernel above reads a *strided* reduced dimension non-coalesced, which silently corrupts
+# large reductions (e.g. `sum(A; dims=2)`); this path avoids that pattern entirely.
+function coalesced_mapreduce_device(f, op, neutral, Rreduce, Rother, R, As...)
+    iother = (get_group_id() - 1) * get_local_size() + get_local_id()
+    gstride = get_num_groups() * get_local_size()
+    @inbounds while iother <= length(Rother)
+        Iother = Rother[iother]
+        Iout = CartesianIndex(Tuple(Iother)..., 1)
+        neut = neutral === nothing ? R[Iout] : neutral
+        val = op(neut, neut)
+        for ireduce in 1:length(Rreduce)
+            Ireduce = Rreduce[ireduce]
+            J = max(Iother, Ireduce)
+            val = op(val, f(_map_getindex(As, J)...))
+        end
+        R[Iout] = val
+        iother += gstride
+    end
+    return
+end
+
 ## COV_EXCL_STOP
 
+# Aurora LTS workaround: the NEO/IGC LTS stack miscompiles *strided* (non-coalesced) global
+# reads inside the reduction kernel, silently corrupting results whenever an input is read
+# along a non-contiguous axis (e.g. `a == transpose(b)`, `sum(transpose(x))`, `ishermitian`).
+# Elementwise copies are NOT affected. `_dense_reduce_input` returns false for any input that
+# reads non-contiguous memory (a transposed/permuted/strided view, or a broadcast containing
+# one), so such inputs get materialized to a dense `oneArray` before reducing. The original
+# symptom was `a == transpose(b)` (the `naive_transpose` example) silently returning false.
+@inline _dense_reduce_input(::oneArray) = true
+@inline _dense_reduce_input(x::Base.ReshapedArray) = _dense_reduce_input(parent(x))
+@inline _dense_reduce_input(::AbstractArray) = false   # Transpose/Adjoint/PermutedDims/SubArray/…
+@inline _dense_reduce_input(::Any) = true              # scalars/Refs/tuples carried in a broadcast
+@inline _dense_reduce_input(bc::Broadcast.Broadcasted) = all(_dense_reduce_input, bc.args)
+
 function GPUArrays.mapreducedim!(f::F, op::OP, R::oneWrappedArray{T},
                                  A::Union{AbstractArray,Broadcast.Broadcasted};
                                  init=nothing) where {F, OP, T}
     Base.check_reducedims(R, A)
     length(A) == 0 && return R # isempty(::Broadcasted) iterates
 
+    # Aurora LTS workaround (see `_dense_reduce_input` above): materialize strided inputs to a
+    # dense array first so every global read in the reduction kernel is coalesced.
+    if oneL0.LTS[] && !_dense_reduce_input(A)
+        Acontig = Broadcast.materialize(Broadcast.broadcasted(f, A))
+        # Only recurse if materialization actually produced a dense `oneArray`. A strided
+        # *device* view (Transpose/PermutedDims of a oneArray, or a broadcast over one)
+        # materializes to a oneArray, but a plain *host* AbstractArray materializes to
+        # another host array that is still not `_dense_reduce_input` — recursing on it would
+        # loop forever (StackOverflowError). Fall through instead so the normal path raises
+        # the usual kernel-conversion error for the unsupported host input.
+        if Acontig isa oneArray
+            return GPUArrays.mapreducedim!(identity, op, R, Acontig; init = init)
+        end
+    end
+
     R_old = R
     # add singleton dimensions to the output container, if needed
     if ndims(R) < ndims(A)
@@ -138,6 +192,43 @@ function GPUArrays.mapreducedim!(f::F, op::OP, R::oneWrappedArray{T},
     # but allows us to write a generalized kernel supporting partial reductions.
     R′ = reshape(R, (size(R)..., 1))
 
+    # Aurora LTS workaround: the workgroup-per-slice kernel below reads a *strided* reduced
+    # dimension non-coalesced, which silently corrupts reductions on this stack (regardless of
+    # output count — it depends on the reduction length, not the number of slices). Whenever
+    # the contiguous leading dimension is NOT reduced (`size(Rreduce, 1) == 1`), use the
+    # coalesced one-work-item-per-slice kernel, whose lanes read consecutive memory. Few-slice
+    # reductions get less parallelism but stay correct; the common many-slice case is also fast.
+    #
+    # `size(Rreduce, 1) == 1` (i.e. dim 1 kept) is the correct predicate, not just a special
+    # case. In `partial_mapreduce_device` adjacent work-item lanes step through adjacent
+    # `Rreduce` entries, whose *fastest* axis is array dim 1. The miscompile bites only when
+    # adjacent lanes land a *full stride* apart on every lane — i.e. when the innermost reduced
+    # axis is itself strided, which is exactly the dim-1-kept case routed to the coalesced
+    # kernel above. When dim 1 is among the reduced dims, adjacent lanes instead read
+    # consecutive memory (stride 1) within each length-n1 run, with only an occasional jump at
+    # a run boundary — not the every-lane full-stride pattern — so the reads stay well-formed
+    # even when a strided dim (e.g. dim 3) is *also* reduced. Verified with exact Int32
+    # reductions (immune to the accumulation-order rounding that makes >2^24-length Float32
+    # sums differ between GPU and CPU): `dims=(1,3)`, `(1,4)` and full `(1,2,3)` reductions at
+    # large strided sizes match the CPU exactly, so these need no extra materialization (see
+    # the "strided mixed reductions" regression test in test/array.jl).
+    if oneL0.LTS[] && size(Rreduce, 1) == 1
+        # cap the group size at what this kernel actually supports on this device (queried
+        # from the driver), rather than a hardcoded 256 that can exceed the kernel's max
+        # work-group size on some devices and fail the launch.
+        cargs = (f, op, init, Rreduce, Rother, R′, A)
+        ckernel = zefunction(
+            coalesced_mapreduce_device,
+            Tuple{Core.Typeof.(kernel_convert.(cargs))...}
+        )
+        items = clamp(length(Rother), 1, launch_configuration(ckernel))
+        groups = min(cld(length(Rother), items), 1024)
+        @oneapi items = items groups = groups coalesced_mapreduce_device(
+            f, op, init, Rreduce, Rother, R′, A
+        )
+        return R_old
+    end
+
     # how many items do we want?
     #
     # items in a group work together to reduce values across the reduction dimensions;
diff --git a/src/oneAPI.jl b/src/oneAPI.jl
index 2abcb64e..7baebfa4 100644
--- a/src/oneAPI.jl
+++ b/src/oneAPI.jl
@@ -17,7 +17,13 @@ using LLVM
 using LLVM.Interop
 using Core: LLVMPtr
 
-using SPIRV_LLVM_Backend_jll, SPIRV_Tools_jll
+import Libdl
+
+# Load both SPIR-V codegen back-ends: GPUCompiler resolves the tool from the target's
+# `backend` field (`:khronos` -> translator, `:llvm` -> LLVM back-end) via a LazyModule that
+# looks the JLL up in `Base.loaded_modules`, so both must be loaded for either path to work.
+# The LTS stack uses the translator; the rolling stack (ONEAPI_LTS=0) uses the LLVM back-end.
+using SPIRV_LLVM_Translator_jll, SPIRV_LLVM_Backend_jll, SPIRV_Tools_jll
 using oneAPI_Support_jll
 
 export oneL0
@@ -142,6 +148,26 @@ function __init__()
             if oneL0.NEO_jll.is_available()
                 # ensure that the OpenCL loader finds the ICD files from our artifacts
                 ENV["OCL_ICD_FILENAMES"] = oneL0.NEO_jll.libigdrcl
+
+                # libsycl (loaded later for oneMKL) carries its own bundled Level Zero loader
+                # that rediscovers the NEO driver by dlopen'ing it by soname rather than
+                # reusing our JLL-loaded module. dlopen NEO here by full path so it is already
+                # resident in this process before libsycl loads: a later
+                # dlopen("libze_intel_gpu.so.1") then resolves to this module by soname without
+                # needing a search path. Required when no system NEO is installed.
+                Libdl.dlopen(oneL0.NEO_jll.libze_intel_gpu; throw_error = false)
+
+                # Extend LD_LIBRARY_PATH with the NEO directory as well. NOTE: this does NOT
+                # affect the running process's own dlopen search — glibc captures
+                # LD_LIBRARY_PATH once at startup — so it serves only *child* processes (e.g.
+                # the multi-worker test suite) that inherit the environment and do their own
+                # path-based driver discovery. In-process discovery is handled by the dlopen
+                # above (and by ZE_ENABLE_ALT_DRIVERS, set in oneL0's __init__).
+                neo_libdir = dirname(oneL0.NEO_jll.libze_intel_gpu)
+                ld = get(ENV, "LD_LIBRARY_PATH", "")
+                if !occursin(neo_libdir, ld)
+                    ENV["LD_LIBRARY_PATH"] = isempty(ld) ? neo_libdir : "$neo_libdir:$ld"
+                end
             end
         end
 
diff --git a/src/oneAPIKernels.jl b/src/oneAPIKernels.jl
index bc6f3218..0a4b70c8 100644
--- a/src/oneAPIKernels.jl
+++ b/src/oneAPIKernels.jl
@@ -264,6 +264,12 @@ function KA.priority!(::oneAPIBackend, prio::Symbol)
         priority = priority_enum
     )
 
+    # Register the replacement queue so `synchronize_all_queues`/`release` can drain it
+    # before freeing a buffer whose in-flight work it references; otherwise all work after
+    # `priority!` runs on an unregistered queue and a freed buffer can be reused while its
+    # kernel is still running (use-after-free → banned context on the LTS NEO stack).
+    oneAPI.register_queue!(ctx, dev, new_queue)
+
     task_local_storage((:ZeCommandQueue, ctx, dev), new_queue)
 
     return nothing
diff --git a/src/pool.jl b/src/pool.jl
index fcc6a3a4..d67f58c8 100644
--- a/src/pool.jl
+++ b/src/pool.jl
@@ -110,6 +110,18 @@ function release(buf::oneL0.AbstractBuffer)
     #    evict(ctx, dev, buf)
     #end
 
+    # NEO (at least the 25.18 LTS release) does not honor the BLOCKING_FREE/DEFER_FREE
+    # policies of zeMemFreeExt: it advertises ZE_extension_memory_free_policies but
+    # unmaps the allocation immediately, even with work in flight that references it.
+    # That turns a GC-driven free of a dead array whose last kernel/copy hasn't retired
+    # into a GPU pagefault, which gets the kernel context banned and makes every later
+    # submission fail with ZE_RESULT_ERROR_UNKNOWN. Synchronize the queues that could
+    # reference this buffer before freeing. (No-op on the rolling stack, which honors
+    # BLOCKING_FREE.)
+    if oneL0.LTS[]
+        synchronize_all_queues(oneL0.context(buf), oneL0.device(buf))
+    end
+
     free(buf; policy=oneL0.ZE_DRIVER_MEMORY_FREE_POLICY_EXT_FLAG_BLOCKING_FREE)
 
     # TODO: queue-ordered free from non-finalizer tasks once we have
diff --git a/src/utils.jl b/src/utils.jl
index f8076a79..ddf4b1f9 100644
--- a/src/utils.jl
+++ b/src/utils.jl
@@ -3,7 +3,8 @@ function versioninfo(io::IO=stdout)
     if Sys.islinux()
         println(io, "Binary dependencies:")
         for jll in [oneL0.NEO_jll, oneL0.NEO_jll.libigc_jll, oneL0.NEO_jll.gmmlib_jll,
-                    SPIRV_LLVM_Backend_jll, SPIRV_Tools_jll, oneAPI_Support_jll]
+                SPIRV_LLVM_Backend_jll, SPIRV_LLVM_Translator_jll, SPIRV_Tools_jll, oneAPI_Support_jll,
+            ]
             name = string(jll)
             print(io, "- $(name[1:end-4]): $(Base.pkgversion(jll))")
             if jll.host_platform !== nothing
diff --git a/test/array.jl b/test/array.jl
index 53eaba96..5d92d759 100644
--- a/test/array.jl
+++ b/test/array.jl
@@ -140,6 +140,27 @@ end
   @test length(b) == 1
 end
 
+@testset "strided mixed reductions" begin
+    # The Aurora LTS IGC miscompiles a reduction kernel's global reads when the *innermost*
+    # reduced axis is strided (dim 1 kept, e.g. `dims=2`); mapreducedim! routes those to a
+    # coalesced kernel. Reductions that also reduce dim 1 (e.g. `dims=(1,3)`) keep a contiguous
+    # innermost axis and stay correct on the workgroup-per-slice kernel — including with a small
+    # leading dim, where the contiguous run is short. Use Int32 (exact, associative) so the
+    # check is immune to Float32 accumulation-order rounding.
+    for (sz, dts) in (
+            ((2, 512, 64), ((1, 3), (2,), (3,), (2, 3), (1, 2, 3), 1)),
+            ((3, 256, 48), ((1, 3), (2,), (1, 2, 3))),
+            ((7, 300, 20), ((1, 3), (2,), (3,))),
+            ((2, 128, 8, 32), ((1, 3), (1, 4), (2, 4), (1, 2, 4))),
+        )
+        A = rand(Int32(1):Int32(4), sz...)
+        dA = oneArray(A)
+        for dt in dts
+            @test Array(sum(dA; dims = dt)) == sum(A; dims = dt)
+        end
+    end
+end
+
 @testset "mapreducedim! returning same type" begin
   R = transpose(oneAPI.zeros(Float32, 2, 3))
   A = oneArray(rand(Float32, 3, 2, 10))
diff --git a/test/level-zero.jl b/test/level-zero.jl
index ed7b2838..bbf8f465 100644
--- a/test/level-zero.jl
+++ b/test/level-zero.jl
@@ -210,14 +210,21 @@ pool = ZeEventPool(ctx, 2)
 signal_event = pool[1]
 wait_event = pool[2]
 
-execute!(queue) do list
-    append_launch!(list, kernel, 1, signal_event, wait_event)
-end
-@test !Base.isdone(signal_event)
+# This is a submit-then-signal pattern: the kernel is gated on `wait_event`, which is
+# only signaled *after* submission. The ONEAPI_SYNC_EACH_SUBMISSION=1 workaround (Aurora
+# LTS) makes `execute!` block in zeCommandQueueSynchronize right after submitting, which
+# would deadlock here since the kernel cannot retire before `wait_event` is signaled. No
+# production code path submits event-gated work, so disable the workaround just here.
+oneL0.sync_each_submission(false) do
+    execute!(queue) do list
+        append_launch!(list, kernel, 1, signal_event, wait_event)
+    end
+    @test !Base.isdone(signal_event)
 
-signal(wait_event)
-synchronize(queue)
-@test Base.isdone(signal_event)
+    signal(wait_event)
+    synchronize(queue)
+    @test Base.isdone(signal_event)
+end
 
 end
 
diff --git a/test/runtests.jl b/test/runtests.jl
index 37bf165d..911b976d 100644
--- a/test/runtests.jl
+++ b/test/runtests.jl
@@ -94,13 +94,16 @@ init_worker_code = quote
         append!(eltypes, [Float64, ComplexF64])
     end
     @static if isdefined(Core, :BFloat16)
-        # Exercising BFloat16 as a generic element type needs the testsuite's CPU
-        # reference path to work, which requires more host-side BFloat16 support than
-        # BFloat16s.jl currently implements (e.g. the div/rem family throws "rem not
-        # defined for BFloat16"). Probe the operations the testsuite relies on, so the
-        # element type enables itself automatically once BFloat16s.jl catches up.
-        # Device-side BFloat16 (oneMKL gemm/gemv, the bfloat16.jl example) is tested
-        # regardless.
+        # Exercising BFloat16 as a generic element type needs (1) hardware support,
+        # (2) a non-LTS stack — the Aurora LTS SPIR-V stack (Khronos translator +
+        # NEO/IGC) cannot codegen native `bfloat` in generic kernels: any kernel that
+        # keeps a bfloat value (e.g. clamp!) fails with InvalidIRError, and declaring
+        # SPV_KHR_bfloat16 crashes the LTS runtime — and (3) a working CPU reference
+        # path, which requires more host-side BFloat16 support than BFloat16s.jl
+        # currently implements (e.g. the div/rem family throws "rem not defined for
+        # BFloat16"). Probe the operations the testsuite relies on, so the element
+        # type enables itself automatically once BFloat16s.jl catches up. Device-side
+        # BFloat16 (oneMKL gemm/gemv, the bfloat16.jl example) is tested regardless.
         bfloat16_host_support() = try
             x, y = BFloat16(3), BFloat16(2)
             mod(x, y); fld(x, y); rand(BFloat16)
@@ -109,7 +112,7 @@ init_worker_code = quote
             false
         end
         const bfloat16_supported = oneAPI._device_supports_bfloat16() &&
-            bfloat16_host_support()
+            !oneL0.LTS[] && bfloat16_host_support()
         if bfloat16_supported
             push!(eltypes, Core.BFloat16)
         end