Intel GPU Architecture: Xe Cores, XMX, Ponte Vecchio and oneAPI SYCL

An Intel GPU is built from Xe-cores. Each Xe-core holds vector engines that run SIMD threads and, on most parts, as many XMX matrix engines, sharing one L1 cache and shared local memory. Xe-cores group into slices and stacks, and you program them through oneAPI: SYCL kernels compiled by DPC++ and run on Level Zero.

That covers four generations that differ a lot up close. Xe-LP in 2020 laptops used execution units with seven threads each. Xe-HPG (Arc A-series) put 16 vector and 16 XMX engines in each Xe-core. Xe-HPC (Ponte Vecchio) doubled the vector width, quadrupled the matrix width and spread 128 Xe-cores over 47 tiles. Xe2 (Battlemage, Lunar Lake) brought the wider design to consumer parts. This guide covers each with Intel's numbers, then SYCL's mapping onto the hardware, worked occupancy and roofline maths, a SYCL kernel, and CUDA migration.

Key takeaways

• The building block is the Xe-core. Xe-HPG has 16 vector engines (XVEs) and 16 XMX engines per Xe-core; Xe-HPC and Xe2 have 8 of each, but wider. Each vector engine runs 8 hardware threads.
• Registers are large and per thread. On Ponte Vecchio each vector engine has 64 KB of registers: 128 registers of 64 bytes per thread in small GRF mode, or 256 in large GRF mode, which halves occupancy.
• A sub-group is one hardware thread. Work-items map to SIMD lanes, a sub-group to one vector engine thread, and a work-group to one Xe-core.
• Ponte Vecchio is a package, not a die: over 100 billion transistors in 47 tiles on five process nodes. The Max 1550 has 128 Xe-cores, 128 GB of HBM2e at 3,276.8 GB/s and 408 MB of L2.
• Peaks need care. From Intel's per-clock figures at 1.6 GHz, the Max 1550 works out to about 52.4 TFLOPS FP32 and 839 TFLOPS dense BF16. These are theoretical peaks, not measurements.
• Software is open and layered. SYCL 2020 is a Khronos standard; DPC++ compiles it and Level Zero runs it. SYCLomatic migrates CUDA, with manual finishing.
• The software around the GPU is a supply chain. Intel published graphics, Level Zero and oneAPI security advisories through 2025 and 2026.

A note on sources and on the code

Specifications come from Intel's oneAPI GPU Optimization Guide (2024.2 and 2025.2), Intel ARK, Intel newsroom releases, the Khronos SYCL 2020 specification, the Level Zero specification, and IEEE Spectrum's report on Intel's ISSCC 2022 Ponte Vecchio paper. Where Intel has not published a figure, we say so. Every computed number was checked in Python.

The code follows the Khronos specification and Intel's documentation but was not compiled or run, because this article was written on a Mac without an Intel GPU. Test it on your own hardware. There are no benchmark results here: every throughput figure is a theoretical peak computed from published per-clock rates and clocks.

If you want the wider context first, our explainers on AI chip architecture and what an AI accelerator is cover the general ideas. How a GPU is made covers packaging and HBM in more depth.

The Xe family at a glance

Intel uses one architecture name, Xe, for a family of designs tuned for different markets. The suffix tells you the target: LP for low power integrated graphics, HPG for high performance gaming, HPC for high performance computing. Xe2 is the second generation, and Intel's optimization guide notes that later generations follow the same pattern, so the third generation is Xe3.

Timeline

Comparison of the Xe-core across generations

This table puts the numbers from Intel's optimization guide side by side. The 2024.2 edition describes Xe-LP, Xe-HPG and Xe-HPC in detail; the 2025.2 edition adds Xe2 columns to its tables.

Two cautions. Intel's guide is not fully consistent: its Xe-slice section says 16 Xe-cores per slice, while its Xe-HPG text says up to 4, which matches the A770's 32 cores in 8 render slices. And the Xe2 XMX width is not in the guide, so we do not quote one.

Names changed along the way

If you read older Intel material you will meet terms that have since been renamed. Knowing the mapping saves confusion.

Intel's guide defines DPAS as "Dot Product and Accumulate Systolic", the elementary operation the XMX engines perform. You will see DPAS in compiler output and in profiler instruction mixes.

Xe-LP: the execution unit design

Xe-LP is worth understanding because it is the bridge between Intel's older integrated graphics and the Xe-core design. Its building block is still the execution unit, and Intel's guide describes it in detail using the Tiger Lake configuration.

Each EU runs 7 hardware threads with simultaneous multithreading. Each thread has 128 general registers of 32 bytes, so 4 KB of register file per thread and 28 KB per EU. The EU has an 8-wide ALU for FP and integer work (SIMD8) and a 2-wide unit for extended math such as square roots and transcendentals (SIMD2). Two EUs can pair to execute SIMD16 instructions.

Per clock, per EU, the guide lists 8 FP32 operations, 16 FP16, 8 INT32, 16 INT16 and 32 INT8 through the DP4A dot product instruction.

Xe-LP dual subslice (one "Xe-core" in later naming)

+---------------------------------------------------------------+
|  16 EUs                                                       |
|  +------+ +------+ +------+ +------+   ... 16 in total         |
|  | EU   | | EU   | | EU   | | EU   |                          |
|  | 7 thr| | 7 thr| | 7 thr| | 7 thr|   each thread: 128 x 32 B |
|  | SIMD8| | SIMD8| | SIMD8| | SIMD8|   registers = 4 KB        |
|  +------+ +------+ +------+ +------+                          |
|                                                               |
|  128 KB shared local memory     data port 128 B/clk           |
+---------------------------------------------------------------+
        6 dual subslices = 96 EUs = 1 slice
        slice: up to 16 MB L2 (3.84 MB last level cache on Tiger Lake config)

The slice in the Tiger Lake configuration has 6 dual subslices, so 96 EUs, which matches the 96 execution units ARK lists for the Core i7-1185G7. The guide's configuration table lists 672 threads (96 times 7) and 1,536 FP32 FLOPs per clock, counting a multiply add as two operations. At the 1.35 GHz maximum graphics clock that ARK lists for the i7-1185G7, that is:

96 EUs x 8 lanes x 2 FLOPs (FMA) = 1,536 FLOPs per clock
1,536 x 1.35e9 = 2.07 TFLOPS FP32 theoretical peak

That figure is a computed ceiling. Real kernels on integrated graphics are usually limited by shared system memory bandwidth long before they reach it.

Xe-HPG: Arc Alchemist and the first Xe-core

Xe-HPG reorganised the hardware around the Xe-core. According to Intel's optimization guide, each Xe-HPG Xe-core contains 16 vector engines and 16 matrix engines.

• Vector engine: 256 bits wide, executing 8 FP32 SIMD operations per clock with fused multiply add. Across the Xe-core that is 256 FP32 operations per clock.
• Matrix engine: 1024 bits wide. Across the Xe-core, the guide lists 4,096 INT8 operations per clock and 2,048 FP16 or BF16 operations per clock.
• Load/store: 512 bytes per clock per Xe-core.

Xe-HPG Xe-core (Arc A770: 32 of these)

+-----------------------------------------------------------------+
|  16 x Xe Vector Engine (256-bit)    16 x XMX engine (1024-bit)  |
|  +-----+-----+-----+-----+          +-----+-----+-----+-----+   |
|  | XVE | XVE | XVE | XVE |          | XMX | XMX | XMX | XMX |   |
|  +-----+-----+-----+-----+          +-----+-----+-----+-----+   |
|  | XVE | XVE | XVE | XVE |          | XMX | XMX | XMX | XMX |   |
|  +-----+-----+-----+-----+   x4     +-----+-----+-----+-----+ x4|
|  each XVE: 8 hardware threads                                   |
|                                                                 |
|  L1 cache / shared local memory      load/store 512 B/clk       |
+-----------------------------------------------------------------+

The A770 has 32 Xe-cores in 8 render slices, so 512 vector engines, 512 XMX engines and 4,096 hardware threads. ARK lists a 2,100 MHz graphics clock, 16 GB of GDDR6 on a 256-bit interface, 560 GB/s of bandwidth, a 225 W board power and TSMC N6. The guide lists 16 MB of last level cache for this configuration.

Worked numbers for the A770

Using the guide's per-clock rates and the ARK graphics clock:

FP32 vector: 32 Xe-cores x 256 FLOPs/clk x 2.1e9      = 17.2 TFLOPS
FP16 XMX:    32 Xe-cores x 2,048 ops/clk x 2.1e9      = 137.6 TFLOPS
INT8 XMX:    32 Xe-cores x 4,096 ops/clk x 2.1e9      = 275.3 TOPS
Memory:      560 GB/s over 256 bits                   = 17.5 Gbps per pin (derived)

The 17.5 Gbps per pin figure is derived by dividing bandwidth by bus width; ARK lists the bandwidth and the width, not the pin rate. All three throughput numbers are theoretical peaks at the listed graphics clock. The card may run faster or slower depending on power and thermals, and Intel's ARK figure is a graphics clock rather than a guaranteed sustained compute clock.

The ratio of compute to bandwidth is the ridge point of a roofline. For FP32 vector work it is 17.2 TFLOPS divided by 560 GB/s, about 30.7 FLOPs per byte. For FP16 on XMX it is about 246 FLOPs per byte. Any kernel that does fewer operations per byte of DRAM traffic than those numbers is bandwidth bound on this card, no matter how well it uses the engines.

Xe-HPC: inside Ponte Vecchio

Ponte Vecchio is the code name for the Intel Data Center GPU Max Series, built for the Aurora supercomputer at Argonne National Laboratory. Intel introduced it on 9 November 2022 as its highest density processor: over 100 billion transistors in a 47 tile package with up to 128 GB of HBM. The release said Aurora would have more than 10,000 blades, each with six Max GPUs and two Xeon Max CPUs; that was the 2022 plan.

The package

Intel presented the design at ISSCC 2022 in a paper titled "Ponte Vecchio: A Multi-Tile 3D Stacked Processor for Exascale Computing", led by Wilfred Gomes. IEEE Spectrum's report on that presentation, published 25 February 2022, gives the structure:

• The package is two mirror image halves joined by Co-EMIB bridges.
• Each half has one base tile, the largest chiplet, which carries cache and the network that links compute tiles to memory.
• Each half has 8 compute tiles built on TSMC N5, and 4 RAMBO SRAM cache tiles built on Intel 7.
• Each half also has 8 blank thermal tiles whose only job is to move heat.
• Xe Link I/O tiles are built on TSMC N7. HBM stacks use their own DRAM process.
• Total silicon area is about 3,100 square millimetres in a 2,330 square millimetre footprint.
• Foveros vertical die to die connections are 36 micrometres apart.
• The design power is 600 W; the package takes 1.8 V and steps it down to about 0.7 V for the compute tiles.

Ponte Vecchio, one package (simplified, not to scale)

   HBM2e HBM2e HBM2e HBM2e            HBM2e HBM2e HBM2e HBM2e
     |     |     |     |                |     |     |     |
 +---------------------------+  EMIB  +---------------------------+
 | Stack 0                   |========| Stack 1                   |
 |  [C][C][C][C] [R][R]      |        |      [R][R] [C][C][C][C]  |
 |  [C][C][C][C] [R][R]      |        |      [R][R] [C][C][C][C]  |
 |  --- Foveros base tile ---|        |--- Foveros base tile ---  |
 |  Xe Link tile             |        |             Xe Link tile  |
 +---------------------------+        +---------------------------+

 C = compute tile (Xe-cores), R = RAMBO cache tile
 Base tile: L2 fabric, memory controllers, I/O; compute and RAMBO tiles
 sit on top of it through Foveros 3D stacking.

Intel counts 47 functional tiles. Secondary sources that quote 63 are adding the 16 thermal tiles, which are silicon but not logic.

Stacks, slices and Xe-cores

To software, each half is an Xe-stack. Intel's guide gives the generic Xe-stack as up to 4 slices, 64 Xe-cores, 64 ray tracing units, 4 hardware contexts, 4 HBM2e controllers, 1 media engine, 8 Xe Links and a shared L2. The two-stack Max 1550 therefore has 8 slices, 128 Xe-cores, 128 ray tracing units, 8 HBM2e controllers and 16 Xe Links.

Max 1550 hierarchy

GPU (2 stacks)
 +-- Xe-stack (x2)
      +-- shared L2 (408 MB total across both stacks)
      +-- 4 HBM2e controllers
      +-- 8 Xe Links
      +-- Xe-slice (x4)
           +-- Xe-core (x16)
                +-- 8 x Xe Vector Engine (512-bit, 8 threads each)
                +-- 8 x XMX engine (4096-bit)
                +-- 512 KB L1 / SLM (SLM up to 128 KB)
                +-- load/store 1024 B/clk

The Xe-HPC Xe-core

The Xe-HPC Xe-core has half as many engines as Xe-HPG, but each one is wider:

• Vector engine: 512 bits, executing 16 FP32 SIMD operations per clock with fused multiply add. The guide lists per Xe-core rates of 512 FP16, 256 FP32 and 128 FP64 operations per clock.
• Matrix engine: 4096 bits. Per Xe-core, 8,192 INT8 and 4,096 FP16 or BF16 operations per clock.
• L1/SLM: 512 KB per Xe-core, with SLM capped at 128 KB per Xe-core and per work-group in the occupancy tables.
• Load/store: 1,024 bytes per clock.

Multiplying out, 128 Xe-cores times 512 KB is 64 MB of L1, which matches the "64MB of L1 cache" in Intel's November 2022 release. The register files add up to the same figure: 1,024 vector engines times 64 KB is another 64 MB.

The product: Max 1550

Intel ARK lists the Max 1550 as follows: 128 Xe-cores, 1,024 XMX engines, 1,024 vector engines, 128 ray tracing units, a 900 MHz base clock and a 1,600 MHz maximum dynamic clock, a 600 W TDP, 128 GB of HBM2e, 3,276.8 GB/s of bandwidth, PCIe Gen 5 x16, and an Xe Link maximum frequency of 53 Gbps. The release adds 408 MB of L2. The Max 1100 is a 300 W PCIe card with 56 Xe-cores and 48 GB of HBM2e; the release also announced a 450 W Max 1350 with 112 Xe-cores and 96 GB.

ARK lists the memory interface width as 1024-bit. Read that as per HBM2e controller rather than for the whole package. With 8 controllers, 3,276.8 GB/s divided by 8 is 409.6 GB/s each, which is exactly what a 1024-bit interface at 3.2 Gbps per pin delivers (1024 times 3.2 divided by 8). That arithmetic is ours; ARK does not spell it out.

Worked peaks for the Max 1550, and a discrepancy

The guide's per-stack table lists 16,384 FP32 FLOPs per clock and 16,384 FP64 FLOPs per clock (counting a multiply add as two), and 262,144 FP16 operations per clock through DP4AS. Using two stacks and the 1.6 GHz maximum clock:

FP32 vector:  2 stacks x 16,384 FLOPs/clk x 1.6e9   = 52.4 TFLOPS
FP64 vector:  2 stacks x 16,384 FLOPs/clk x 1.6e9   = 52.4 TFLOPS (per table)
BF16 XMX:     128 Xe-cores x 4,096 ops/clk x 1.6e9  = 838.9 TFLOPS
INT8 XMX:     128 Xe-cores x 8,192 ops/clk x 1.6e9  = 1,677.7 TOPS

The FP64 line needs a caveat. The same guide's prose says each Xe-core does 128 FP64 operations per clock, which multiplies out to 26.2 TFLOPS, half the table figure. We could not resolve the difference from Intel's pages, and ARK lists no TFLOPS figure. IEEE Spectrum's February 2022 report gave more than 45 TFLOPS of FP32, a pre-launch number.

All of these are theoretical peaks at the maximum dynamic clock, which a 600 W part does not hold under every load.

The roofline for the Max 1550

                 838.9 TFLOPS (BF16 XMX peak) ----------------------
                /
               /  ridge at 256 FLOP/byte
              /
     52.4 TFLOPS (FP32 vector peak) ------------
            /   ridge at 16 FLOP/byte
           /
          /  slope = 3,276.8 GB/s (HBM2e)
         /
   -----+-------------+------------------+----------> arithmetic intensity
       0.17           16                256          (FLOPs per byte of HBM traffic)
     (SAXPY)

Two worked kernels show where real code lands.

SAXPY in FP32 (y = a*x + y) does 2 FLOPs per element and moves 12 bytes: read x, read y, write y. That is 0.167 FLOPs per byte. At 3,276.8 GB/s the ceiling is about 546 GFLOPS, around 1 percent of the FP32 vector peak. Nothing about the vector engines matters for this kernel. Only bandwidth does.

A 4096 by 4096 by 4096 GEMM in BF16 with FP32 output does 2N cubed, about 1.37 times 10 to the 11th FLOPs. If each matrix is read or written once, traffic is N squared times (2 + 2 + 4) bytes, about 134 MB. The intensity is 1,024 FLOPs per byte, well past the 256 ridge, so the kernel is compute bound and could in principle finish in about 0.16 ms at the XMX peak. Real GEMMs read tiles more than once, so effective intensity is lower, but it stays far above the ridge for large sizes. This is why matrix libraries spend so much effort on tiling into L1/SLM and registers: the XMX engines can only be kept fed if each byte fetched from HBM is reused many times.

For a sense of what memory bound means in time, a vector add of 2 to the 26th FP32 elements touches about 805 MB (two reads and a write). At the full 3,276.8 GB/s that is about 246 microseconds, a floor no kernel tuning can beat.

Xe Link

Each stack has 8 Xe Links, 16 per Max 1550, and ARK gives the Xe Link maximum frequency as 53 Gbps. Intel's 2022 release describes a four GPU OAM carrier board joined by Xe Link. We found no Intel page stating lanes per link, so we do not compute an aggregate bandwidth.

Xe2: Battlemage and Lunar Lake

Xe2 is the second generation. It appears as Xe2-LPG integrated graphics, for example in Core Ultra 200V (Lunar Lake) laptops (Intel's guide uses the Core Ultra 7 265 as its Xe2-LPG example), and as Xe2-HPG in Arc B-series discrete cards.

From the 2025.2 optimization guide, Xe2 looks structurally closer to Xe-HPC than to Xe-HPG:

• 8 vector engines per Xe-core, down from 16 on Xe-HPG.
• 512-bit registers, up from 256 bits on Xe-HPG.
• 128 or 256 GRF registers per thread, small or large register mode, the same scheme as Xe-HPC.
• Sub-group sizes of 16 and 32 only; SIMD8 is gone.
• Native double precision listed as yes.
• 128 KB of SLM per Xe-core, and up to 128 KB per work-group, double the Xe-HPG limit.
• L1 per Xe-core of 256 KB on the B580 and 192 KB on the Core Ultra 7 265.
• Last level cache of 18 MB on the B580 and 8 MB on the Core Ultra 7 265.

Arc B580

Intel ARK lists the Arc B580 with 20 Xe2 cores, 160 XMX engines, 5 render slices, 20 ray tracing units, a 2,670 MHz graphics clock, 12 GB of GDDR6 on a 192-bit interface, 456 GB/s of bandwidth, a 190 W board power, PCIe 4.0 x8, TSMC N5, and 233 peak INT8 TOPS.

The 160 XMX engines across 20 cores give 8 XMX engines per Xe2 core. The bandwidth works out to 19 Gbps per pin (456 times 8 divided by 192), again derived rather than listed.

An arithmetic check shows why per engine rates should not be reverse engineered from marketing peaks. Dividing 233 INT8 TOPS by 160 XMX engines and 2.67 GHz gives about 545 operations per clock per engine, not a power of two. At 512 per clock, the listed clock would give about 218.7 TOPS; 233 TOPS would need about 2.84 GHz. The TOPS figure probably assumes a different clock, but Intel does not say so and the guide does not publish the Xe2 per engine rate. We treat 233 TOPS as Intel's stated peak and derive nothing from it.

Lunar Lake

Intel introduced Core Ultra 200V on 3 September 2024. Its release says the built-in GPU includes as many as eight second generation Xe-cores, eight ray tracing units and XMX engines with up to 67 TOPS, and that the platform reaches up to 120 total TOPS. Intel's earlier Lunar Lake architecture fact sheet described the Xe2 GPU as delivering more than 60 TOPS, with an editor's note that its Xe2 claims were updated on 4 June 2024. The two figures are consistent: "more than 60" was rounded before launch, and 67 is the launch figure. The 67 TOPS is a peak.

After Xe2: Xe3 and Xe3P

Intel's guide notes that the third generation is called Xe3. On the data center side, Intel's OCP 2026 page describes the Intel GPU D-series, formerly code-named Crescent Island, as built on the Xe3P architecture and designed for AI inference. It states up to 480 GB of LPDDR5X memory, 1.5 TB/s of memory bandwidth and 1.3 petaflops of FP8 compute. The page does not say whether that compute figure is peak or sustained, and it gives no Xe-core count, XMX count, power figure or ship date. We have not found an Intel optimization guide entry for Xe3P, so this post goes no deeper on it. Treat anything else you read about Xe3P internals as unconfirmed until Intel's developer documentation covers it.

Inside the Xe-core: threads, registers and SIMD

Now the part that matters most when you write kernels: how a vector engine actually executes code.

Hardware threads

A vector engine is a multithreaded SIMD processor. On Xe-HPG, Xe-HPC and Xe2, each vector engine holds 8 hardware threads, each with its own instruction pointer and its own slice of the register file. Every clock, the engine's thread arbiter can pick ready instructions from different threads and issue them to the execution pipes. When one thread waits on memory, others keep the pipes busy. This is the same latency hiding idea GPUs have always used, but with fewer, fatter threads than a design that issues 32-wide warps.

One Xe-HPC vector engine over time (8 threads, small GRF mode)

clock:   1    2    3    4    5    6    7    8    9   10
T0:     FMA  FMA  LD...........................wait   FMA
T1:          FMA  FMA  LD......................wait
T2:               FMA  FMA  FMA  LD............
T3:                         FMA  FMA  FMA  FMA
T4:                                   FMA  FMA ...
...
Pipes stay busy as long as some thread has a ready instruction.

The diagram is schematic, not a cycle accurate trace; Intel does not publish per cycle issue rules at that level in the optimization guide.

The general register file

Each hardware thread owns a dedicated block of the general register file (GRF). On Xe-HPC, Intel's guide is precise: each vector engine has 64 KB of registers.

Both rows multiply to 64 KB per engine. The guide states the trade-off plainly: using a large register file brings Xe-core occupancy down by 50 percent. In return, a kernel with heavy register pressure stops spilling to memory. The guide's own experiment reports that small GRF mode performed better when loop unrolling was off and large GRF mode performed better when unrolling was on, which is exactly the trade you would expect: unrolling raises register pressure.

On Ponte Vecchio you choose the mode at compile time:

# Ahead-of-time compile for PVC, force large GRF for every kernel
icpx -fsycl -fsycl-targets=spir64_gen \
  -ftarget-register-alloc-mode=pvc:large \
  -Xsycl-target-backend "-device pvc" test.cpp

# Just-in-time compile, let the compiler choose per kernel
icpx -fsycl -ftarget-register-alloc-mode=pvc:auto test.cpp

The guide says the only device currently supported by this option is pvc, and that without the option Intel's graphics compiler picks a mode per kernel by heuristics on Linux. The modes are default, small, large and auto.

On Xe-LP the register file was smaller per thread: 128 registers of 32 bytes, 4 KB per thread, 7 threads per EU, 28 KB per EU.

SIMD width and sub-groups

The second key idea is that the compiler, not the hardware, picks how many work-items one hardware thread carries. A kernel compiled at SIMD16 puts 16 work-items in each hardware thread; at SIMD32, 32 work-items. SYCL exposes this as the sub-group size. Intel's 2025.2 guide states the mapping directly: each work-item maps to a SIMD lane, a sub-group maps to a vector engine hardware thread, and a work-group runs on one Xe-core.

On Xe-HPC, a 512-bit register holds 16 FP32 values. A SIMD16 FP32 operation fills one register per operand; a SIMD32 FP32 operation spans two. SIMD32 doubles the work-items per thread, which can help hide latency with fewer threads, but it doubles the registers each variable consumes, which can force spills or large GRF mode.

You can pin the sub-group size per kernel with the SYCL 2020 attribute:

q.parallel_for(sycl::nd_range<1>(global, local),
  [=](sycl::nd_item<1> it) [[sycl::reqd_sub_group_size(16)]] {
    // this kernel is compiled at SIMD16
  });

Pick a size the device supports. On Xe-HPC and Xe2 that means 16 or 32; asking for 8 there is an error.

XMX engines and DPAS

The XMX engine is a systolic array. Data flows through a grid of multiply accumulate cells, and each instruction computes a small matrix product and adds it to an accumulator. The instruction is DPAS.

Width tells you throughput. On Xe-HPG, each 1024-bit XMX engine contributes 128 FP16 operations per clock (2,048 per Xe-core divided by 16 engines). On Xe-HPC, each 4096-bit engine contributes 512 FP16 or BF16 operations per clock (4,096 divided by 8). Four times the width, four times the per engine throughput, with half as many engines per core.

The matrix shapes a single sub-group can hand to the engine come from the joint_matrix tests in Intel's guide:

One BF16 tile operation on PVC at 8 by 16 by 16 is 2 times 8 times 16 times 16, or 4,096 FLOPs. The N of 16 matches the SIMD16 sub-group: each work-item in the sub-group holds one column's share of the result.

Programming XMX with joint_matrix

In SYCL you reach XMX through the sycl::ext::oneapi::experimental::matrix extension. The key rule from Intel's guide is that every work-item in the sub-group must call the joint matrix functions together, with no divergence. The core loop looks like this, adapted from the guide's example:

using namespace sycl::ext::oneapi::experimental::matrix;

// TM x TN tile of C per sub-group; K walked in steps of TK
sycl::sub_group sg = item.get_sub_group();

joint_matrix<sycl::sub_group, sycl::ext::oneapi::bfloat16, use::a, TM, TK, layout::row_major> a;
joint_matrix<sycl::sub_group, sycl::ext::oneapi::bfloat16, use::b, TK, TN, layout::row_major> b;
joint_matrix<sycl::sub_group, float, use::accumulator, TM, TN> c;

joint_matrix_fill(sg, c, 0.0f);
for (size_t k = 0; k < K / TK; ++k) {
  joint_matrix_load(sg, a, pA + (row * TM) * K + k * TK, K);
  joint_matrix_load(sg, b, pB + (k * TK) * N + col * TN, N);
  joint_matrix_mad(sg, c, a, b, c);          // lowers to DPAS on XMX
}
joint_matrix_apply(sg, c, [=](float &x) { x *= alpha; });
joint_matrix_store(sg, c, pC + (row * TM) * N + col * TN, N, layout::row_major);

Two warnings from the guide. First, there is no emulation: the device must have XMX hardware, and you can check by running sycl-ls --verbose and looking for the ext_intel_matrix aspect. Second, the API has changed between releases; older 2023 guides used a return value form of joint_matrix_mad. Match the code to the compiler version you actually ship.

For production GEMMs, most teams call oneMKL (now developed as oneMath under the UXL Foundation for the open interfaces) or oneDNN rather than writing joint_matrix by hand. The libraries carry hand tuned tiling for each target.

The memory hierarchy

Per work-item / per thread      GRF registers        8 KB or 16 KB per thread (PVC)
        |
Per Xe-core                     L1 + SLM             512 KB combined (PVC); SLM up to 128 KB
        |
Per stack                       L2 (RAMBO + base)    408 MB total across 2 stacks (PVC)
        |
Per stack                       HBM2e, 4 controllers 64 GB and about 1.64 TB/s per stack (derived)
        |
Between stacks / GPUs           Xe Link, PCIe Gen 5

The per stack HBM figures are half the package figures, 128 GB and 3,276.8 GB/s, on the assumption that memory is split evenly between the two stacks, which matches 4 HBM2e controllers per stack in the guide.

Shared local memory

SLM is the equivalent of CUDA shared memory: a scratchpad visible to all work-items in a work-group, allocated per work-group. The limits from the 2025.2 guide:

Allocation is rounded. The guide notes that SLM requests round up to 2, 4, 8, 16, 24, 32, 48, 64, 96 or 128 KB, so a 65 KB request occupies 96 KB, which leaves room for only one such work-group per Xe-core.

L2 and the last level cache

The L2 on Ponte Vecchio is enormous by GPU standards: 408 MB per package. Part of it lives in the RAMBO tiles stacked on the base tile. For a working set that fits, L2 behaves like a bandwidth amplifier: data is fetched from HBM once and reused many times. If your kernel's working set is a few hundred megabytes, it is worth measuring whether reordering loops to stay inside L2 pays off. Intel does not publish an L2 bandwidth figure on the pages cited here, so we cannot put an L2 ceiling on the roofline.

On Xe-HPG the guide lists 16 MB of last level cache for the A770, and on Xe2 18 MB for the B580 and 8 MB for the Core Ultra 7 265.

Occupancy, worked on the Max 1550

Occupancy on Intel GPUs is measured in hardware threads, not work-items. Intel's guide defines Xe-core occupancy as the ratio of hardware threads in use in an Xe-core to the total available, and GPU occupancy the same way across the whole GPU.

The Max 1550 numbers, straight from the guide: 8 threads per vector engine, 8 vector engines per Xe-core, so 64 hardware threads per Xe-core, and 128 Xe-cores, so 8,192 hardware threads.

Step 1: threads per work-group

A work-group of W work-items compiled at sub-group size S needs W divided by S hardware threads. At sub-group 32, a 1024 work-item group needs 32 threads; a 256 work-item group needs 8.

Step 2: work-groups per Xe-core

Divide the Xe-core's thread budget by threads per work-group, then check SLM.

The last two rows come from the guide's own table. With 24 threads per group, two groups use 48 of 64 threads (75 percent) and a third would need 72. With 12 threads per group, five groups use 60 of 64 (93.75 percent).

The guide also lists two launch failures worth remembering. A work-group of 896 with that global range fails because the work-groups are non-uniform, and a work-group of 1,152 fails because it exceeds the 1,024 maximum.

Step 3: the SLM limit

Now add SLM. At work-group size 1024 and sub-group 32, two groups fit by threads. Each can use up to 64 KB of SLM and both still fit in the 128 KB per Xe-core. At 65 KB per group, rounding takes each to 96 KB, only one group fits, and occupancy drops to 50 percent. At more than 128 KB, the guide reports the kernel fails to launch with ERROR_OUT_OF_RESOURCES.

Step 4: large GRF mode

Large GRF mode halves threads per vector engine to 4, so an Xe-core has 32 threads. A 1024 work-item group at sub-group 32 needs all 32, so only one group fits. A 256 work-item group at sub-group 16 needs 16 threads, so two fit. Always recompute occupancy when you flip the register mode.

Step 5: waves

The guide's worked example runs a 64 by 120 by 128 problem with work-groups of 1 by 8 by 128 (1,024 work-items). That is 960 work-groups. At 2 groups per Xe-core and 128 Xe-cores, one wave holds 256 groups, so the launch takes 960 divided by 256, or 3.75 waves: three full waves and a last wave at 75 percent. Shapes that leave a nearly empty final wave waste a slice of runtime, which matters for short kernels.

Here is the same calculation as a small script you can adapt. It was run with Python 3 while writing this post and printed the numbers above.

def xe_core_occupancy(wg, sg, threads_per_core=64, slm_per_wg_kb=0, slm_core_kb=128):
    rounding = [0, 2, 4, 8, 16, 24, 32, 48, 64, 96, 128]
    slm = next((r for r in rounding if r >= slm_per_wg_kb), None)
    if slm is None or wg > 1024 or wg % sg:
        return "launch fails"
    t = wg // sg
    by_threads = threads_per_core // t
    by_slm = slm_core_kb // slm if slm else by_threads
    groups = min(by_threads, by_slm)
    return groups, groups * t / threads_per_core

print(xe_core_occupancy(1024, 32, slm_per_wg_kb=65))   # (1, 0.5)
print(xe_core_occupancy(384, 32))                      # (5, 0.9375)
print(xe_core_occupancy(1024, 32, threads_per_core=32))  # large GRF: (1, 1.0)

When occupancy does not matter

High occupancy is a means, not a goal. A GEMM in large GRF mode at 50 percent occupancy can still be faster, because each thread keeps more of its tile in registers and issues DPAS back to back. A bandwidth bound kernel needs just enough threads to saturate memory. Use occupancy to diagnose stalls, not as a score.

The software stack: oneAPI, DPC++, IGC and Level Zero

Intel's GPU software is layered. From top to bottom:

+--------------------------------------------------------------+
| Frameworks and apps: PyTorch, TensorFlow, HPC codes          |
+--------------------------------------------------------------+
| Libraries: oneMKL / oneMath, oneDNN, oneCCL, oneDPL          |
+--------------------------------------------------------------+
| Language: SYCL 2020 (Khronos) + Intel extensions, OpenMP     |
|           offload, OpenCL                                    |
+--------------------------------------------------------------+
| Compiler: Intel oneAPI DPC++/C++ (icpx), based on LLVM       |
|           emits SPIR-V for JIT or native code for AOT        |
+--------------------------------------------------------------+
| SYCL runtime with backends: Level Zero, OpenCL               |
+--------------------------------------------------------------+
| Level Zero loader + Intel compute runtime (GPU driver)       |
| Intel Graphics Compiler (IGC) turns SPIR-V into GPU ISA      |
+--------------------------------------------------------------+
| Kernel mode driver (i915 or xe on Linux), firmware           |
+--------------------------------------------------------------+
| Hardware: Xe-cores, XMX, L1/SLM, L2, HBM or GDDR             |
+--------------------------------------------------------------+

oneAPI and the UXL Foundation

oneAPI is both a specification and Intel's toolkit. The specification now sits with the Unified Acceleration (UXL) Foundation, a vendor neutral Joint Development Foundation project whose site lists oneDPL, oneDNN, oneCCL, oneDAL, oneTBB and oneMath. Intel's oneAPI Base Toolkit packages the compiler, runtime and Intel's builds of these libraries.

SYCL 2020

SYCL is a Khronos standard for single source C++ heterogeneous programming. The version fetched for this article is the SYCL 2020 specification, revision 12. Two defaults from the specification are worth knowing on day one: queues are out of order by default and order work by dependencies, and you can request an in-order queue with a property to get first in, first out behaviour much like a CUDA stream. Unified shared memory (USM) gives you malloc_device, malloc_host and malloc_shared, which map closely to device, pinned host and managed memory.

DPC++ and the compiler flow

Intel's compiler driver is icpx. With -fsycl it compiles the host code natively and the device code to SPIR-V. At run time, the SYCL runtime hands the SPIR-V to the Intel Graphics Compiler through the driver, which compiles it to the GPU's native instruction set. That is just-in-time compilation and it costs time on first launch. Ahead-of-time compilation with -fsycl-targets=spir64_gen and a device name such as pvc does that step at build time instead, and it is also how you pass device specific options like the GRF mode.

Level Zero

Level Zero is the lowest public layer. Its specification (version 1.20.41 at the time of writing) describes the API's aim as providing "direct-to-metal interfaces to offload accelerator devices". It is designed for runtimes and libraries rather than most applications, exposes drivers, devices, contexts, command queues, command lists, events, modules and kernels, and makes the application responsible for releasing every handle, since the API neither reference counts nor garbage collects.

A minimal device discovery in C looks like this. It follows the specification's API names and was not run for this post.

#include <level_zero/ze_api.h>
#include <stdio.h>

int main(void) {
    if (zeInit(ZE_INIT_FLAG_GPU_ONLY) != ZE_RESULT_SUCCESS) return 1;

    uint32_t ndrv = 0;
    zeDriverGet(&ndrv, NULL);
    ze_driver_handle_t drv[8];
    if (ndrv > 8) ndrv = 8;
    zeDriverGet(&ndrv, drv);

    for (uint32_t d = 0; d < ndrv; ++d) {
        uint32_t ndev = 0;
        zeDeviceGet(drv[d], &ndev, NULL);
        ze_device_handle_t dev[16];
        if (ndev > 16) ndev = 16;
        zeDeviceGet(drv[d], &ndev, dev);
        for (uint32_t i = 0; i < ndev; ++i) {
            ze_device_properties_t p = { .stype = ZE_STRUCTURE_TYPE_DEVICE_PROPERTIES };
            zeDeviceGetProperties(dev[i], &p);
            printf("%s: %u slices, %u subslices/slice, %u EUs/subslice, %u threads/EU\n",
                   p.name, p.numSlices, p.numSubslicesPerSlice,
                   p.numEUsPerSubslice, p.numThreadsPerEU);
        }
    }
    return 0;
}

The property names still use the older EU and subslice terms. How a two stack device is reported depends on driver settings, so check rather than assume. For a quicker look, sycl-ls --verbose lists devices and aspects such as ext_intel_matrix, and ONEAPI_DEVICE_SELECTOR=level_zero:gpu restricts a program to the Level Zero GPU backend. Those two answer most "why is my kernel running on the CPU" questions.

A SYCL kernel, worked through

Here is a complete program in three steps: a vector add with USM, then a reduction that uses SLM and sub-group operations, with the hardware mapping explained at each line. It follows the SYCL 2020 specification and was not compiled for this article.

Step 1: vector add with USM

#include <sycl/sycl.hpp>
#include <vector>
#include <iostream>

int main() {
  constexpr size_t N = 1 << 26;                       // 67,108,864 floats
  sycl::queue q{sycl::gpu_selector_v, sycl::property::queue::in_order{}};
  std::cout << q.get_device().get_info<sycl::info::device::name>() << "\n";

  float *a = sycl::malloc_device<float>(N, q);
  float *b = sycl::malloc_device<float>(N, q);
  float *c = sycl::malloc_device<float>(N, q);

  std::vector<float> ha(N, 1.0f), hb(N, 2.0f), hc(N);
  q.memcpy(a, ha.data(), N * sizeof(float));
  q.memcpy(b, hb.data(), N * sizeof(float));

  q.parallel_for(sycl::range<1>(N), [=](sycl::id<1> i) {
    c[i] = a[i] + b[i];
  });

  q.memcpy(hc.data(), c, N * sizeof(float)).wait();
  std::cout << "c[0] = " << hc[0] << "\n";            // expect 3

  sycl::free(a, q); sycl::free(b, q); sycl::free(c, q);
}

What happens on the hardware:

• sycl::range<1>(N) lets the runtime choose the work-group size. The compiler picks a SIMD width, typically 16 or 32 on Xe-HPC.
• Each sub-group becomes one hardware thread on one vector engine. At SIMD16, 67,108,864 work-items become 4,194,304 hardware thread instances, scheduled across 8,192 thread slots on a Max 1550.
• Each work-item loads two floats and stores one. The kernel moves about 805 MB and, as computed earlier, cannot finish faster than about 246 microseconds at the 3,276.8 GB/s peak. That is a floor, not a measurement.
• The in-order queue means the copies and the kernel run in submission order, the way a single CUDA stream would.

Step 2: a reduction with SLM and sub-groups

A sum reduction exercises every level of the hierarchy: registers, sub-group shuffles, SLM and global atomics.

float sum_device(sycl::queue &q, const float *x, size_t n) {
  constexpr size_t WG = 256;          // work-items per work-group
  constexpr size_t SG = 16;           // sub-group size, valid on Xe-HPC and Xe2
  size_t groups = (n + WG - 1) / WG;

  float *out = sycl::malloc_shared<float>(1, q);
  *out = 0.0f;

  q.submit([&](sycl::handler &h) {
    // one float per sub-group: 256 / 16 = 16 slots
    sycl::local_accessor<float, 1> partial(sycl::range<1>(WG / SG), h);

    h.parallel_for(sycl::nd_range<1>(groups * WG, WG),
      [=](sycl::nd_item<1> it) [[sycl::reqd_sub_group_size(SG)]] {
        size_t gid = it.get_global_id(0);
        float v = gid < n ? x[gid] : 0.0f;

        // 1. reduce within the sub-group: lanes of one hardware thread
        sycl::sub_group sg = it.get_sub_group();
        v = sycl::reduce_over_group(sg, v, sycl::plus<float>());

        // 2. one lane per sub-group writes to SLM
        if (sg.leader()) partial[sg.get_group_id()[0]] = v;
        sycl::group_barrier(it.get_group());

        // 3. first sub-group reduces the 16 partials
        if (sg.get_group_id()[0] == 0) {
          float p = sg.get_local_id()[0] < WG / SG ? partial[sg.get_local_id()[0]] : 0.0f;
          p = sycl::reduce_over_group(sg, p, sycl::plus<float>());
          if (sg.leader()) {
            sycl::atomic_ref<float, sycl::memory_order::relaxed,
                             sycl::memory_scope::device,
                             sycl::access::address_space::global_space> acc(*out);
            acc.fetch_add(p);
          }
        }
      });
  }).wait();

  float r = *out;
  sycl::free(out, q);
  return r;
}

The mapping, line by line:

• Work-group of 256 at sub-group 16 needs 16 hardware threads. On a Max 1550 Xe-core with 64 threads, 4 groups fit by threads. SLM use is 16 floats, 64 bytes, which rounds to the 2 KB minimum, so SLM does not limit occupancy. Occupancy is 100 percent in small GRF mode.
• reduce_over_group on a sub-group stays inside one hardware thread. The lanes exchange values through the register file, with no SLM traffic and no barrier. This is the SYCL counterpart of a CUDA warp shuffle reduction.
• local_accessor allocates SLM per work-group. The 16 partial sums sit in the Xe-core's L1/SLM array.
• group_barrier synchronises the 16 hardware threads of the work-group, the counterpart of __syncthreads().
• atomic_ref with device scope adds each group's result to one float in global memory. With many work-groups that atomic becomes a hot spot; for very large inputs a two pass reduction, or the built-in SYCL reduction object, is usually better.

Floating point addition is not associative, so the result can differ in the last bits from a serial CPU sum and between runs, because atomics complete in a nondeterministic order. Tests should compare with a tolerance.

Step 3: build and run

# Source the oneAPI environment (path depends on your install)
source /opt/intel/oneapi/setvars.sh

# List devices the SYCL runtime can see
sycl-ls

# JIT build: device code as SPIR-V, compiled by the driver at first launch
icpx -fsycl -O3 reduce.cpp -o reduce

# AOT build for Ponte Vecchio
icpx -fsycl -O3 -fsycl-targets=spir64_gen -Xsycl-target-backend "-device pvc" \
  reduce.cpp -o reduce_pvc

# Force the Level Zero GPU backend
ONEAPI_DEVICE_SELECTOR=level_zero:gpu ./reduce

These commands follow Intel's documentation; we did not run them for this post.

Mapping CUDA concepts to SYCL

If you know CUDA, most ideas carry over. The table lists the closest counterparts and the traps.

The dimension order trap catches nearly everyone once. In CUDA, threadIdx.x varies fastest; in a two dimensional SYCL nd_range, dimension 1 does. Keeping x as dimension 0 turns coalesced loads into strided ones.

The other trap is sub-group size. Code that assumes 32 lanes, such as a 32-bit ballot mask or arrays sized per warp, breaks or wastes resources when the compiler picks 16. Query sg.get_local_range() or pin the size.

Migrating CUDA code with SYCLomatic

SYCLomatic is Intel's open source project for migrating CUDA source to SYCL. Its repository describes it as a project to assist developers in migrating existing code, and it is explicit about the limit: "Final code editing and verification is a manual process done by the developer." The tool binary is c2s, and dpct is an alias for it. The repository lists CUDA versions 8.0, 9.x, 10.x, 11.x and 12.0 to 12.9 as supported inputs, and warns that none of its branches are stable or rigorously tested for production quality control, so production teams usually use the released versions.

A typical flow for a Make based project:

# 1. Capture the real compiler invocations into a compilation database
intercept-build make

# 2. Migrate everything in the database into ./out
c2s -p compile_commands.json --in-root=. --out-root=out

# 3. Read the warnings c2s inserted as comments, fix each by hand
grep -rn "DPCT1" out/ | head

# 4. Build with the oneAPI compiler and run your tests
icpx -fsycl out/*.cpp -o app

These steps follow the project's documented usage and were not run for this post. What to expect in the output:

• Kernel launches become parallel_for over an nd_range, with the dimension order swapped as described above.
• Runtime calls like cudaMalloc become USM calls, often through helper headers that ship with the tool. Many teams later remove those helpers to get plain SYCL.
• Library calls to cuBLAS and similar map to oneMKL where a counterpart exists.
• Warnings appear as comments where the tool could not prove a translation correct, for example around warp size assumptions, inline PTX, or error handling semantics that differ.

Inline PTX, hard coded warp size, unusual texture modes and vendor intrinsics need hand work. Budget a test pass and a performance pass, since a kernel tuned for 32-wide warps is rarely optimal for 16-wide sub-groups.

Step by step: tuning a SYCL kernel for an Intel GPU

1. Confirm the device. Run sycl-ls, set ONEAPI_DEVICE_SELECTOR=level_zero:gpu, and print the queue's device name.
2. Compute the roofline first. Below the ridge point (about 16 FLOPs per byte for FP32 on a Max 1550, about 31 on an A770), optimise memory traffic, not arithmetic.
3. Coalesce on the fastest dimension, the last nd_range dimension.
4. Choose sub-group size deliberately. Try 16 and 32 and watch register pressure at 32.
5. Size work-groups for occupancy, using your device's threads per Xe-core and SLM rounding.
6. Decide the GRF mode on Ponte Vecchio: try auto, then small and large, and measure.
7. Use libraries for GEMM and convolution. Write joint_matrix only when no library covers your shape or fusion.
8. Compile ahead of time for production to avoid first launch JIT cost.
9. Profile occupancy, stalls and XMX use, and compare with steps 2 and 5.
10. Pin and record your toolchain (compiler, runtime and driver versions) so results reproduce and advisories can be matched to what you ship.

Common mistakes

• Treating a sub-group like a 32-wide warp. Xe-HPC and Xe2 support 16 and 32 only, Xe-HPG supports 8, 16 and 32, and the compiler picks unless you say otherwise.
• Ignoring the dimension order. A two dimensional CUDA kernel translated without swapping x to the last SYCL dimension produces strided memory access.
• Forgetting SLM rounding. A 65 KB SLM request uses 96 KB and halves the work-groups per Xe-core on a Max 1550.
• Exceeding the work-group limit. The maximum is 1,024 on Xe-HPG, Xe-HPC and Xe2 but 512 on Xe-LP. A kernel that launches on one fails on the other.
• Assuming FP64 everywhere. Intel's guide lists native double precision as unavailable on Xe-HPG. Code that needs FP64 should check the device aspect and fail clearly.
• Quoting peak TFLOPS as performance. Every TFLOPS and TOPS figure in this post is a theoretical peak at a stated clock. Bandwidth bound kernels never approach them.
• Mixing generations in one AOT target. An AOT binary built for pvc will not run on an Arc card. Ship a JIT path or build for each target.
• Expecting joint_matrix to fall back. There is no emulation on devices without XMX; check ext_intel_matrix before choosing that code path.
• Leaving the GPU software stack unpinned. Container images that pull the newest runtime at build time produce different binaries week to week and make advisories hard to match.

FAQ

What is an Xe-core?

An Xe-core is the compute building block of Intel GPUs from Xe-HPG onward. It contains vector engines for SIMD work, XMX matrix engines on most products, and a shared L1 cache and shared local memory. A SYCL work-group always runs on a single Xe-core.

What is the difference between Xe-LP, Xe-HPG, Xe-HPC and Xe2?

Xe-LP is the 2020 integrated graphics design built from execution units with 7 threads each. Xe-HPG (Arc A-series) uses Xe-cores with 16 vector and 16 XMX engines. Xe-HPC (Ponte Vecchio) uses Xe-cores with 8 wider vector and 8 wider XMX engines, HBM2e and a multi-tile package. Xe2 (Battlemage, Lunar Lake) brings 8 vector engines per Xe-core and 512-bit registers to consumer products.

What are XMX engines?

XMX engines are Intel's matrix units, systolic arrays that execute DPAS (dot product accumulate systolic) instructions. On Xe-HPC each Xe-core does 4,096 FP16 or BF16 operations per clock on XMX, according to Intel's guide. You use them through oneMKL, oneDNN, frameworks, or the SYCL joint_matrix extension.

How many Xe-cores does Ponte Vecchio have?

The Data Center GPU Max 1550 has 128 Xe-cores across two stacks of 64, with 1,024 vector engines and 1,024 XMX engines, according to Intel ARK. The Max 1100 has 56.

What is the sub-group size on Intel GPUs?

It depends on device and kernel. Intel's guide lists 16 and 32 for Xe-HPC and Xe2, and 8, 16 and 32 for Xe-LP and Xe-HPG. One sub-group runs as one hardware thread on a vector engine.

Can SYCLomatic convert all CUDA code automatically?

No. The project states that final code editing and verification is manual. It handles most runtime and kernel syntax, then flags what it cannot translate with warnings for a developer to resolve.

What does large GRF mode do?

On Ponte Vecchio, large GRF mode gives each hardware thread 256 registers of 64 bytes instead of 128, which halves the threads per vector engine from 8 to 4. It cuts register spills for heavy kernels at the cost of 50 percent lower occupancy.

Is the Data Center GPU Max still sold?

Intel ARK lists the Max 1550 launch as Q1 2023 and its expected discontinuance as January 2026. Check ARK and your vendor for current availability and support terms before planning around it.

The GPU software supply chain

Everything above the silicon in the stack diagram is software, and software has vulnerabilities. Intel's security center lists, among others, Intel Graphics advisories on 12 May 2026 (INTEL-SA-01402) and 10 February 2026 (INTEL-SA-01385), Graphics Software advisories on 11 November 2025 and 12 August 2025, a oneAPI Math Kernel Library advisory on 11 November 2025 (INTEL-SA-01366), a oneAPI Toolkit advisory on 12 August 2025 (INTEL-SA-01285), a oneAPI Level Zero advisory and a oneAPI DPC++/C++ Compiler advisory on 13 May 2025 (INTEL-SA-01274 and INTEL-SA-01243), and a Data Center GPU Max Series advisory on 13 August 2024 (INTEL-SA-01094).

Two things make this harder than ordinary dependency management. The layers are versioned together, so patching the runtime can mean retesting the driver and framework. And the pieces arrive by different routes: OS packages, vendor repositories, toolkit installers, base images and Python wheels. Teams often know their framework version but not which Level Zero loader or oneDNN build is inside the image.

The fix is the same as for any other software: keep a software bill of materials for every image and host, record exact versions of the runtime and libraries, match those against advisories, and rebuild when a fix lands. If you self host models on accelerators, our post on GPU cost per million tokens covers the economics; the inventory discipline applies to whichever accelerator you pick.

Where Safeguard fits

Safeguard does not design, program or tune GPUs, and it does not make SYCL or oneAPI code run faster. It has no role in Intel's drivers or firmware, and it does not scan GPU firmware or kernel drivers on a host. The fit is the software layer that sits in your repositories and container images.

Safeguard scans repositories, container images and SBOM files (safeguard scan --dir ., --image <name>, --sbom <file>) and matches the components it finds against vulnerability databases, with reachability analysis to rank what your code actually calls. If your inference image bundles oneDNN, oneMKL or a framework's Intel GPU support layer, those components and their versions show up in the inventory, and the platform can show which projects contain a given component and version when an Intel advisory lands. Policies and gates evaluated in the platform can stop a build that carries a known vulnerable version. Host level drivers and firmware outside your images are still your OS and vendor update process.

A rebuilt GPU image pulls its runtime, libraries and framework layers again, and every one of them is open source or vendor software with its own release schedule. That is why the gap between a published fix and a fix in production matters as much for an accelerator stack as for a web service, and why the practices below still apply.

More in this series

Other posts in our GPUs, CPUs, QPUs and lithography series:

• AMD GPU Architecture Explained: GCN, RDNA, CDNA, Infinity Fabric and HIP
• CPU vs GPU vs QPU Explained: How Each Works and Why They Coexist
• CUDA Programming Model Explained: Grids, Blocks, Warps and Memory
• GPU Memory and Interconnects: HBM, GDDR, NVLink, PCIe, CXL
• CUDA vs HIP vs SYCL vs OpenCL vs Metal vs Vulkan vs Triton Compared
• GPU Software Supply Chain: Drivers, CUDA, cuDNN, NCCL and Your SBOM
• How a QPU Works: Qubit Chips, Fridges, Control Pulses and Readout
• How CPUs Work: Pipelines, Out-of-Order, Caches, Branches, SIMD
• How EUV Lithography Machines Work: Tin Plasma, Mirrors and High-NA
• How GPUs Work: Graphics Pipeline, Shader Cores, Ray Tracing and Compute
• Inside an NVIDIA SM: Warp Schedulers, Registers, Tensor Cores
• NVIDIA GPU Architecture History: From Tesla and Fermi to Blackwell
• NVIDIA vs AMD vs Intel vs Apple GPUs: Architectures Compared
• Optimizing a CUDA Matrix Multiply Step by Step, From Naive to Tensor Cores
• Photolithography Explained: Contact Printing to DUV Immersion and Overlay
• Tile-Based GPUs Explained: Apple, Mali, Adreno, PowerVR and TBDR

Fix it without waiting for upstream

Whatever the vulnerability, most of the risk sits in two gaps. The first is the time between a fix being published and that fix actually running in your production builds. The second, and worse, is the time before any fix exists at all, while a maintainer works out a patch and a release schedule. Waiting through both gaps is a habit, not a requirement. Safeguard gives teams three ways to close them.

1. Autonomous remediation: pull requests, not tickets

When a fixed version exists, Safeguard's Griffin models open a tested, policy gated pull request that moves the affected dependency to the safe version. You can run it on one repository or in bulk across every project that carries the vulnerable version. Remediation covers eight categories, including dependencies, licenses, security posture and findings. Reachability analysis ranks the vulnerable code paths your application actually calls first, so the first pull requests you merge are the ones that matter most. Autonomous remediation ships switched off, with every category set to Manual, so your team decides how much runs on its own and what still needs a human approval.

2. Gold Open Source: the safe version in seconds, free

gold.safeguard.sh is free and needs no login. Search any CVE or package to see affected and fixed versions, CVSS, EPSS, CISA KEV status and a recommended safe version, or use the "Am I affected?" checker. The same data is available through a free read API, README status badges and RSS threat feeds. The safeguard-sh/gold-scan-action GitHub Action can fail a build on a known exploited vulnerability, on any known vulnerability, or never, so a version you have just removed cannot quietly come back in next week's merge.

3. Gold Registry: start from components that are already fixed

registry.safeguard.sh provides hardened, drop in replacements for upstream packages and container images. Every Gold artifact is held continuously to zero exploitable critical or high CVEs, zero malware and a completed license check, with SLSA L3 or higher provenance, a Sigstore signature and a public evidence page you can hand to an auditor. An artifact loses its Gold badge if it regresses. Teams can point their own package directory at the registry so hardened equivalents resolve without changing how developers install anything.

When upstream has no fix yet, or will not ship one, Safeguard first works to get the fix accepted upstream. If the issue is severe or upstream will not move, Safeguard builds a fork, fixes the vulnerability there and publishes a hardened Gold build, so you are not stuck on a maintainer's release schedule. You can also request a hardened, compatibility tested build of a specific package. This fork and patch flow is in early access today, so ask which packages it covers before you plan around it.

One honest limit: when the vulnerable code lives in a vendor appliance, a hosted service or an operating system runtime you do not build yourself, the fix still has to come from that vendor or distribution. The steps above apply to the open source dependencies and container images you build and ship, which is where most teams carry the largest and least tracked share of their exposure.

See how Safeguard would help

The first demo is a real Safeguard CLI session. You install it, run safeguard to open the interactive terminal and ask it to fix the lodash CVE in a repository. It reads the finding and advisory, asks before any command that changes something, edits package.json to lodash 4.17.21, runs npm install and npm test with your approval, prints a summary and pushes nothing. The same flow applies to the Python and C++ dependencies around a SYCL or PyTorch project, and /scan . gives you a first inventory of what the repository pulls in.

Click through it at your own pace, or open it full screen.

The second demo sets up the local runner on your own machine, approves it with its device code, then scans from the terminal. GPU build hosts and inference images often live on private networks, and this is how a scan can run where the images are built.

Click through it at your own pace, or open it full screen.

The third demo connects Claude to the Safeguard MCP server with OAuth and shows how you choose which tools Claude can use. Once connected, you can ask which of your projects carry a given library version, a quick way to see which GPU images need a rebuild after an advisory.

Click through it at your own pace, or open it full screen.

To try it on your own projects, download Safeguard or read the CLI documentation.

Quick check

Sources

• Intel oneAPI GPU Optimization Guide 2024.2: Intel Xe GPU Architecture
• Intel oneAPI GPU Optimization Guide 2025.2: Intel Xe GPU Architecture
• Intel oneAPI GPU Optimization Guide 2025.2: Thread Mapping and GPU Occupancy
• Intel oneAPI GPU Optimization Guide 2025.2: Small Register Mode vs. Large Register Mode
• Intel oneAPI GPU Optimization Guide 2025.2: Programming Intel XMX Using SYCL Joint Matrix Extension
• Intel ARK: Intel Data Center GPU Max 1550 specifications
• Intel ARK: Intel Arc A770 Graphics 16GB specifications
• Intel ARK: Intel Arc B580 Graphics specifications
• Intel ARK: Intel Core i7-1185G7 processor specifications
• Intel newsroom: Intel Max Series Brings Breakthrough Memory Bandwidth and Performance to HPC and AI (9 November 2022)
• Intel newsroom: New Core Ultra Processors Deliver Breakthrough Performance, Efficiency for the AI PC Age (3 September 2024)
• Intel fact sheet: Intel Unveils Lunar Lake Architecture
• Intel newsroom: Intel at OCP 2026
• IEEE Spectrum: Behind Intel's HPC Chip that Will Pierce the Exascale Barrier (25 February 2022)
• Khronos: SYCL 2020 Specification, revision 12
• oneAPI Level Zero Specification
• SYCLomatic repository
• UXL Foundation
• Intel Product Security Center advisories

Last reviewed: 25 October 2026. Specifications change between driver and documentation releases; check the linked Intel pages for the latest figures.