CUDA vs HIP vs SYCL vs OpenCL vs Metal vs Vulkan vs Triton Compared

CUDA, HIP, SYCL, OpenCL, Metal, Vulkan compute and Triton all express the same idea: launch a grid of small thread groups that share fast on-chip memory and synchronise at barriers. They differ in who can run the code, which libraries come with it, and how much tuning the compiler does for you.

CUDA runs only on NVIDIA GPUs and has the deepest library and tool ecosystem. HIP is AMD's near copy of the CUDA API for AMD GPUs. SYCL and OpenCL are Khronos standards that several vendors implement. Metal is Apple's API for Apple GPUs. Vulkan compute is the compute half of a cross-vendor graphics API. Triton is a Python tile language that a compiler lowers to NVIDIA and AMD GPUs. This guide writes one kernel, a tiled matrix multiply, in all seven, maps the vocabulary between them, and walks through the published evidence on which ones keep their performance when they move between vendors.

Key takeaways

• Every model in this guide exposes the same three level hierarchy: a grid (or ND range, or dispatch) of groups, each group a set of threads that can share on-chip memory and meet at a barrier. Learn it once and the rest is mostly vocabulary.
• The hardware's SIMD width leaks through. NVIDIA warps are 32 threads; AMD documents 64 thread wavefronts on CDNA and 32 on RDNA; the Apple M5 GPU used for this article reports a SIMD width of 32. Code that hard codes 32 breaks quietly on a 64 wide machine.
• Group size limits differ. A 16 by 16 group (256 threads) is legal on every CUDA GPU, but Vulkan core only guarantees 128 invocations per workgroup, and Vulkan 1.4 raises that floor to 256. Always query.
• Tiling is what makes a matrix multiply fast in every model. For a square FP32 tile of side T, arithmetic intensity against device memory is T/4 FLOP per byte, so 16 by 16 tiles give 4 and 32 by 32 tiles give 8, up from 0.25 for the naive kernel.
• Published studies show portable models can come close to native ones. A 2024 study across NVIDIA and AMD supercomputers found SYCL "most consistently" performance portable for its tests, followed by RAJA and Kokkos. A 2023 study found SYCL on GPUs slightly ahead of native CUDA and HIP on average for bandwidth bound codes. Neither result transfers automatically to your workload.
• Triton and NVIDIA's newer CUDA Tile model move the hardest tuning work (shared memory staging, tensor core use) from the programmer to the compiler. You trade fine control for productivity.
• Whichever model you choose, you inherit a software stack: drivers, toolkits, runtime libraries, compilers and container images. That stack has its own versions and vulnerabilities and belongs in your SBOM.

A note on the code: the Metal and OpenCL kernels in this article were compiled and run on an Apple M5 Mac and checked against a CPU result (maximum absolute error 0.0 on a 100 by 100 integer valued test). The CUDA, HIP, SYCL, Vulkan GLSL and Triton kernels follow the vendor and standards documentation but were not compiled or run, because this Mac has no NVIDIA, AMD or Intel GPU and no Vulkan shader compiler installed. No timing numbers in this article come from our own runs; every performance claim is attributed to a published study.

The one kernel we will write seven times

A matrix multiply, C = A × B, is the right test piece. It is simple enough to fit on a screen, it is the core of most AI workloads, and the gap between a naive version and a tuned one is enormous, which makes it a good probe of how much each model helps you.

We use square N by N FP32 matrices in row major order. Element (i, j) of A lives at A[i * N + j].

The naive kernel and why it is slow

The naive kernel gives each thread one output element and has it walk a full row of A and a full column of B:

// One thread per output element; every operand comes from device memory.
float acc = 0.0f;
for (int k = 0; k < N; ++k)
    acc += A[row * N + k] * B[k * N + col];
C[row * N + col] = acc;

Each iteration performs one fused multiply add, which counts as 2 floating point operations, and loads two 4 byte floats. If nothing is reused from cache, that is 2 FLOP for 8 bytes, an arithmetic intensity of 0.25 FLOP per byte. For N = 4096 the whole product is 2N³ = 137,438,953,472 FLOP, about 137.4 GFLOP, and without reuse the kernel would request about 549.8 GB from device memory. Caches recover some of that in practice, but you cannot count on them to do the programmer's job.

Tiling: the idea every model shares

Tiling splits C into T by T tiles. One group of T × T threads owns one output tile. It walks along the K dimension in steps of T, and at each step the group cooperatively copies one T by T tile of A and one of B into on-chip memory, waits at a barrier, multiplies the two tiles out of on-chip memory, waits again, and moves on.

              K dimension, walked in steps of T
          +-------+-------+-------+-------+
   A      | A0    | A1    | A2    | A3    |   row band of A owned by this group
          +-------+-------+-------+-------+

          +-------+
   B      | B0    |   column band of B owned by this group
          +-------+
          | B1    |
          +-------+
          | B2    |
          +-------+
          | B3    |
          +-------+

   step t: copy At and Bt into on-chip memory (one element per thread)
           barrier
           every thread does T multiply adds from on-chip memory
           barrier
   after the last step each thread writes one element of its C tile

Per step, the group loads 2T² floats (8T² bytes) from device memory and performs 2T³ FLOP. The intensity against device memory is therefore 2T³ / 8T² = T/4 FLOP per byte. Here are the worked numbers for N = 4096, checked in python:

The last row is a trap. A 64 by 64 group with one thread per element would need 4,096 threads, which no model in this guide allows in one group. The CUDA Programming Guide says a thread block "may contain up to 1024 threads" on current GPUs, and Metal's Feature Set Tables give the same 1,024 for Apple4 and later. Real kernels get to larger tiles by having each thread compute several outputs, which is called register blocking. We keep T = 16 and one output per thread in all seven versions so they stay comparable.

What the intensity buys you: a roofline sketch

The roofline model says attainable throughput is the lower of the compute peak and intensity times memory bandwidth. To show the shape without quoting a real part, take a hypothetical GPU with a 20 TFLOP/s FP32 peak and 2 TB/s of memory bandwidth. These are round numbers for illustration, not any vendor's specification.

The ridge point of this imaginary device is 20 / 2 = 10 FLOP per byte. Below it you are memory bound; above it, compute bound. These are ceilings, not predictions. Real kernels lose throughput to barriers, address arithmetic, occupancy limits and bank conflicts, which is exactly the work that libraries and compilers like Triton try to take off your hands. For the hardware side of this argument, our explainer on AI chip architecture and the one on what an AI accelerator is cover why bandwidth, not FLOPs, so often sets the limit.

A terminology map across all seven models

Before the code, here is the vocabulary. The concepts line up closely; the names do not.

Three rows deserve a comment.

The SIMD bundle row is where portability leaks. The CUDA Programming Guide states that "threads are organized into groups of 32 threads called warps." AMD's HIP programming model documentation says the warp size "is architecture dependent and always fixed: 64 threads for CDNA architectures, 32 threads for RDNA architectures." SYCL's specification leaves the sub group size implementation defined, and says the size of an individual sub group "is invariant for the duration of a kernel's execution." Metal exposes threads_per_simdgroup as a built in, and the Apple M5 GPU we used reported a threadExecutionWidth of 32 for our kernel.

The shared memory row hides different capacities. Apple's Metal Feature Set Tables list 32 KB of threadgroup memory for Apple4 family GPUs and later, and our M5 reported maxThreadgroupMemoryLength of 32,768 bytes. The Vulkan specification's required limits table guarantees only 16,384 bytes of maxComputeSharedMemorySize on core Vulkan, with 32,768 listed for the Vulkan Roadmap 2026 profile. Our 16 by 16 FP32 kernel needs 2,048 bytes, so it fits everywhere; a heavily register blocked kernel may not.

The Triton row is mostly "not exposed" by design. Triton's programmer works at the level of a whole tile, and the compiler decides how threads, shared memory and tensor cores are used.

A dated history of the seven models

We have not given a first release date for HIP or Metal because we did not find a primary source that states one clearly. The current documents are HIP 7.16.0 documentation on AMD's ROCm site and the Metal Shading Language Specification version 4.1 from Apple.

CUDA: the reference point

CUDA is NVIDIA's platform: a C++ dialect compiled by nvcc, a runtime and driver API, and a large set of libraries. NVIDIA's CUDA Programming Guide describes the model this way: threads are grouped into thread blocks, blocks into a grid, and within a block threads execute in warps of 32 under a single instruction, multiple threads (SIMT) model. From compute capability 9.0, thread block clusters let a group of blocks read and write each other's shared memory, which the guide calls distributed shared memory.

The tiled kernel in CUDA

// CUDA C++: tiled FP32 matrix multiply, one output element per thread.
// Follows the CUDA Programming Guide; not compiled or run for this article.
#define TILE 16

__global__ void matmul_tiled(const float* A, const float* B, float* C, int N)
{
    __shared__ float As[TILE][TILE];
    __shared__ float Bs[TILE][TILE];

    const int row = blockIdx.y * TILE + threadIdx.y;
    const int col = blockIdx.x * TILE + threadIdx.x;
    float acc = 0.0f;

    for (int t = 0; t < (N + TILE - 1) / TILE; ++t) {
        const int aCol = t * TILE + threadIdx.x;
        const int bRow = t * TILE + threadIdx.y;
        As[threadIdx.y][threadIdx.x] = (row < N && aCol < N) ? A[row * N + aCol] : 0.0f;
        Bs[threadIdx.y][threadIdx.x] = (bRow < N && col < N) ? B[bRow * N + col] : 0.0f;
        __syncthreads();

        for (int k = 0; k < TILE; ++k)
            acc += As[threadIdx.y][k] * Bs[k][threadIdx.x];
        __syncthreads();
    }
    if (row < N && col < N)
        C[row * N + col] = acc;
}

// Host side launch
dim3 block(TILE, TILE);
dim3 grid((N + TILE - 1) / TILE, (N + TILE - 1) / TILE);
matmul_tiled<<<grid, block>>>(dA, dB, dC, N);
cudaDeviceSynchronize();

A few details matter. threadIdx.x varies fastest, so consecutive threads in a warp read consecutive addresses of A and B, which lets the hardware coalesce the loads into a few wide transactions. The CUDA guide spells out the cost of getting this wrong with a worked example: if every thread in a warp forces its own 32 byte transaction, the warp moves 32 bytes × 32 threads = 1,024 bytes to use only 128. Inside the inner loop, As[threadIdx.y][k] is the same address for every thread in a row, which the hardware broadcasts, and Bs[k][threadIdx.x] is consecutive, so neither access causes shared memory bank conflicts. The two __syncthreads() calls are both required: the first makes sure the tile is fully loaded before anyone reads it, and the second makes sure nobody overwrites the tile while a slower warp is still reading it.

The guard on loads and stores (row < N, aCol < N) lets N be any size, not just a multiple of 16. Out of range elements are loaded as zero, which does not change the sum.

Beyond the hand written kernel

Production code on NVIDIA hardware rarely uses a hand written kernel like this. cuBLAS provides tuned GEMM routines, CUTLASS provides C++ templates for building custom ones, and the warp level matrix interfaces (WMMA and the PTX MMA instructions) target tensor cores. The CUDA Programming Guide now also documents CUDA Tile, which it describes as "a different approach to writing GPU kernel code than the single instruction multiple thread (SIMT) model." In a tile kernel "the programmer thinks at the level of an entire block," and the guide says tile programming lets the compiler use features such as the tensor memory accelerator and tensor cores. CUDA Tile is available in Python through the cuda.tile package and in C++ "from version 13.3 onward" of the CUDA Toolkit. That is NVIDIA's own move toward the tile style that Triton popularised.

CUDA strengths and costs

Strengths: the broadest library set (cuBLAS, cuDNN, cuFFT, NCCL, TensorRT, CUTLASS), mature profilers (Nsight Systems and Nsight Compute), and documentation that tracks each new architecture.

Costs: it runs on NVIDIA GPUs only. Your binaries and containers carry a dependency on specific driver and toolkit versions, and moving to another vendor means a port.

HIP: CUDA's shape on AMD hardware

HIP is AMD's C++ runtime API and kernel language, part of ROCm. AMD's documentation says GPU programmers with CUDA experience "will find the HIP API straightforward" and that you "can quickly port NVIDIA CUDA applications to run on AMD GPUs." The ROCm libraries mirror the CUDA ones: hipBLAS, hipFFT and others provide API compatibility with their CUDA equivalents.

The tiled kernel in HIP

The kernel body is the same as CUDA's. Only the header and host API names change.

// HIP C++: identical kernel body to the CUDA version.
// Follows the HIP documentation; not compiled or run for this article.
#include <hip/hip_runtime.h>
#define TILE 16

__global__ void matmul_tiled(const float* A, const float* B, float* C, int N)
{
    __shared__ float As[TILE][TILE];   // lives in the LDS (local data share) on AMD GPUs
    __shared__ float Bs[TILE][TILE];
    const int row = blockIdx.y * TILE + threadIdx.y;
    const int col = blockIdx.x * TILE + threadIdx.x;
    float acc = 0.0f;
    for (int t = 0; t < (N + TILE - 1) / TILE; ++t) {
        const int aCol = t * TILE + threadIdx.x;
        const int bRow = t * TILE + threadIdx.y;
        As[threadIdx.y][threadIdx.x] = (row < N && aCol < N) ? A[row * N + aCol] : 0.0f;
        Bs[threadIdx.y][threadIdx.x] = (bRow < N && col < N) ? B[bRow * N + col] : 0.0f;
        __syncthreads();
        for (int k = 0; k < TILE; ++k)
            acc += As[threadIdx.y][k] * Bs[k][threadIdx.x];
        __syncthreads();
    }
    if (row < N && col < N) C[row * N + col] = acc;
}

// Host side
float *dA, *dB, *dC;
hipMalloc(&dA, N * N * sizeof(float));
hipMemcpy(dA, hA, N * N * sizeof(float), hipMemcpyHostToDevice);
// ... same for dB, dC
dim3 block(TILE, TILE);
dim3 grid((N + TILE - 1) / TILE, (N + TILE - 1) / TILE);
hipLaunchKernelGGL(matmul_tiled, grid, block, 0, 0, dA, dB, dC, N);
hipDeviceSynchronize();

With a 16 by 16 block, a CDNA GPU runs 256 / 64 = 4 wavefronts per block, while an NVIDIA GPU runs 256 / 32 = 8 warps. The kernel does not care, because it never assumes a warp size. Many real kernels do.

The warp size trap, worked

AMD's HIP porting guide is direct about it: "Code should not assume a warp size of 32 or 64, as AMD GPU architectures have different warp sizes. The warpSize built-in should be used in device code." The guide also warns about lane masks. A lane mask has one bit per thread in a warp. On a 32 wide warp, a full mask fits in a 32 bit integer, whose largest value is 4,294,967,295. On a 64 wide wavefront you need a 64 bit integer, whose largest value is 18,446,744,073,709,551,615. CUDA code that writes 1u << lane with a 32 bit unsigned silently drops the upper 32 lanes on CDNA. The guide's example computes laneId = threadIdx.x % warpSize and uses that, not a literal 32.

A warp level reduction is where this bites most often:

// Wrong on a 64 wide wavefront: hard coded 32 and a 32 bit mask.
for (int offset = 16; offset > 0; offset /= 2)
    v += __shfl_down_sync(0xffffffffu, v, offset);

// Portable shape: start from warpSize / 2 and use the platform's mask type.
for (int offset = warpSize / 2; offset > 0; offset /= 2)
    v += __shfl_down(v, offset);   // HIP form; check the HIP docs for mask variants

Porting tools

AMD ships HIPIFY in two forms. The documentation describes hipify-clang as a Clang based tool that "parses code, translates it into an Abstract Syntax Tree, and generates the HIP source," and hipify-perl as a text based tool that can also convert code that is not syntactically correct, which makes it "easier to set up and use, but is not as powerful." The porting guide also documents scan scripts that report how much of a directory can be converted automatically, and an in place mode that keeps a .prehip backup of each file.

The guide points out smaller differences too. For example, amdclang++ does not support the --maxregcount option that nvcc accepts, and AMD recommends __launch_bounds__ instead.

HIP strengths and costs

Strengths: a CUDA developer is productive on day one, and HIPIFY handles the mechanical renaming. The ROCm library set covers the common needs.

Costs: it targets AMD GPUs, so a single HIP codebase is not a route to Intel or Apple hardware. Performance tuning still differs: wavefront width, LDS size and matrix instruction shapes are AMD specific. A 2020 paper from the Ginkgo linear algebra project, porting from CUDA to HIP, reported having to add missing features such as cooperative groups to the HIP framework and discussed "the performance price of compiling HIP code for AMD architectures." Things have moved since 2020, which is a reason to test on your own release, not to assume.

SYCL: standard C++ across vendors

SYCL is a Khronos standard for single source heterogeneous programming in standard C++. Khronos announced the SYCL 2020 final specification on 9 February 2021, saying it is "Based on C++17" and builds on SYCL 1.2.1. The current document is revision 12, dated 6 August 2026.

There is no single SYCL compiler. Intel's oneAPI DPC++ compiler is "a LLVM-based compiler project that implements compiler and runtime support for the SYCL language," developed in the open in Intel's LLVM repository. AdaptiveCpp, formerly hipSYCL and Open SYCL, describes itself as "the independent, community-driven modern platform for C++-based heterogeneous programming models targeting CPUs and GPUs from all major vendors." The 2023 Reguly study notes that both compilers then supported GPUs from all three major vendors.

The tiled kernel in SYCL

// SYCL 2020: tiled FP32 matmul using USM pointers and local_accessor.
// Follows the SYCL 2020 specification; not compiled or run for this article.
#include <sycl/sycl.hpp>
constexpr int TILE = 16;

void matmul(sycl::queue& q, const float* A, const float* B, float* C, int N)
{
    const size_t G = ((N + TILE - 1) / TILE) * TILE;   // round global size up

    q.submit([&](sycl::handler& h) {
        sycl::local_accessor<float, 2> As(sycl::range<2>(TILE, TILE), h);
        sycl::local_accessor<float, 2> Bs(sycl::range<2>(TILE, TILE), h);

        h.parallel_for(sycl::nd_range<2>({G, G}, {TILE, TILE}),
                       [=](sycl::nd_item<2> it) {
            const int row = it.get_global_id(0);
            const int col = it.get_global_id(1);   // dimension 1 varies fastest
            const int ly  = it.get_local_id(0);
            const int lx  = it.get_local_id(1);
            float acc = 0.0f;

            for (int t = 0; t < (N + TILE - 1) / TILE; ++t) {
                const int aCol = t * TILE + lx;
                const int bRow = t * TILE + ly;
                As[ly][lx] = (row < N && aCol < N) ? A[row * N + aCol] : 0.0f;
                Bs[ly][lx] = (bRow < N && col < N) ? B[bRow * N + col] : 0.0f;
                sycl::group_barrier(it.get_group());

                for (int k = 0; k < TILE; ++k)
                    acc += As[ly][k] * Bs[k][lx];
                sycl::group_barrier(it.get_group());
            }
            if (row < N && col < N) C[row * N + col] = acc;
        });
    }).wait();
}

// Host side: USM device allocations
// sycl::queue q{sycl::gpu_selector_v};
// float* dA = sycl::malloc_device<float>(N * N, q);
// q.memcpy(dA, hA, N * N * sizeof(float)).wait();

Three things differ from CUDA. First, the group shared memory is declared on the host side of the submission as a local_accessor, not inside the kernel. SYCL 2020 deprecated the older target::local accessor in favour of this class. Second, the index order is reversed from CUDA's habit: in SYCL the last dimension of a range varies fastest, so we put the column in dimension 1 to keep loads coalesced. Getting this backwards compiles fine and runs slowly. Third, the barrier is a free function, group_barrier, that takes the group, and the same function works on a sub group.

Migrating CUDA to SYCL

Intel's open source SYCLomatic project helps here. Its README says it assists "in migrating their existing code written in different programming languages to the SYCL C++ heterogeneous programming model," with a c2s binary for CUDA code, and adds that "Final code editing and verification is a manual process done by the developer." Treat the output as a first draft, not a finished port.

SYCL strengths and costs

Strengths: standard C++, several independent implementations, and one code base that can target NVIDIA, AMD and Intel GPUs plus CPUs. The published evidence on performance portability, covered below, is better for SYCL than many people expect.

Costs: you depend on the SYCL implementation's back end for each vendor, and features that live outside the standard (tensor core access, vendor specific atomics) arrive as extensions that differ by implementation. Library coverage for each vendor comes through oneMKL style interfaces or the vendor's own library underneath, so check that the routines you need exist on every target.

OpenCL: the original portable standard

OpenCL was the first cross-vendor GPU compute standard. Khronos announced the ratification and public release of OpenCL 1.0 at SIGGRAPH Asia in Singapore, with a press release dated 9 December 2008, calling it "the first open, royalty-free standard for cross-platform, parallel programming." OpenCL 3.0, published as a provisional specification on 27 April 2020, made most features beyond OpenCL 1.2 optional so more devices could claim conformance. The current unified specification, v3.1.2 dated 17 September 2026, includes a summary of changes from OpenCL 3.0 to OpenCL 3.1 and lists SPIR-V as a required intermediate language for OpenCL 2.1, 2.2 and 3.1 or newer devices.

The tiled kernel in OpenCL C

This kernel was compiled and run on an Apple M5 through Apple's OpenCL implementation, which reported itself as OpenCL 1.2, and matched a CPU reference exactly on a 100 by 100 integer valued test.

// OpenCL C: tiled FP32 matmul. Compiled and run on Apple M5 (OpenCL 1.2)
// for this article; max absolute error vs CPU reference: 0.0 (N = 100).
#define TILE 16

__kernel void matmul_tiled(__global const float* A,
                           __global const float* B,
                           __global float*       C,
                           const int             N)
{
    __local float As[TILE][TILE];
    __local float Bs[TILE][TILE];

    const int lx  = get_local_id(0);
    const int ly  = get_local_id(1);
    const int col = get_global_id(0);
    const int row = get_global_id(1);
    float acc = 0.0f;

    for (int t = 0; t < (N + TILE - 1) / TILE; ++t) {
        const int aCol = t * TILE + lx;
        const int bRow = t * TILE + ly;
        As[ly][lx] = (row < N && aCol < N) ? A[row * N + aCol] : 0.0f;
        Bs[ly][lx] = (bRow < N && col < N) ? B[bRow * N + col] : 0.0f;
        barrier(CLK_LOCAL_MEM_FENCE);

        for (int k = 0; k < TILE; ++k) {
            acc += As[ly][k] * Bs[k][lx];
        }
        barrier(CLK_LOCAL_MEM_FENCE);
    }
    if (row < N && col < N) {
        C[row * N + col] = acc;
    }
}

The host side is where OpenCL shows its age. You pick a platform and device, create a context and a command queue, build the program from source at run time, create the kernel, set each argument by index, and enqueue it:

size_t G = ((N + 15) / 16) * 16;
size_t global[2] = { G, G };
size_t local[2]  = { 16, 16 };
clSetKernelArg(k, 0, sizeof(cl_mem), &dA);
clSetKernelArg(k, 1, sizeof(cl_mem), &dB);
clSetKernelArg(k, 2, sizeof(cl_mem), &dC);
clSetKernelArg(k, 3, sizeof(int),    &N);
clEnqueueNDRangeKernel(queue, k, 2, NULL, global, local, 0, NULL, NULL);

Notice that OpenCL uses dimension 0 as the fastest varying index, like CUDA's x and unlike SYCL's last dimension. Non-uniform work-groups are an optional device feature in the current specification (reported through CL_DEVICE_NON_UNIFORM_WORK_GROUP_SUPPORT, a query that did not exist before 3.0), so the portable choice is to round the global size up to a multiple of the local size and rely on the bounds checks, as we do with G.

Evidence on CUDA versus OpenCL

Karimi, Dickson and Hamze compared near identical CUDA and OpenCL kernels from a quantum Monte Carlo application in 2010. Their abstract reports that with NVIDIA's tools, "converting a CUDA kernel to an OpenCL kernel involves minimal modifications," while compiling the same kernel with ATI's tools needed more changes.

OpenCL strengths and costs

Strengths: the widest device reach of any model here, including CPUs, GPUs, FPGAs and many embedded parts, with a stable C API that rarely breaks.

Costs: verbose host code, uneven support for optional features across vendors, and a thin AI library ecosystem.

Metal: Apple's GPU API

Metal is Apple's graphics and compute API for its GPUs. Kernels are written in the Metal Shading Language (MSL), a C++ based language; the current specification is version 4.1. Apple's Metal Feature Set Tables list per family limits. For Apple4 family GPUs and later they give a maximum of 1,024 threads per threadgroup and 32 KB of total threadgroup memory; Apple2 and Apple3 allow 512 threads.

The tiled kernel in Metal

We compiled this kernel at run time with MTLDevice.makeLibrary(source:) on an Apple M5, dispatched it from a small Swift host program and compared the result to a CPU reference on a 100 by 100 integer valued test. The maximum absolute error was 0.0. The device reported threadExecutionWidth 32, maxTotalThreadsPerThreadgroup 1,024 and maxThreadgroupMemoryLength 32,768 bytes for this pipeline. We did not time it, and we make no performance claim from it.

// Metal Shading Language: tiled FP32 matmul. Compiled and run on Apple M5
// for this article; max absolute error vs CPU reference: 0.0 (N = 100).
#include <metal_stdlib>
using namespace metal;

#define TILE 16

kernel void matmul_tiled(device const float* A [[buffer(0)]],
                         device const float* B [[buffer(1)]],
                         device float*       C [[buffer(2)]],
                         constant uint&      N [[buffer(3)]],
                         uint2 gid [[thread_position_in_grid]],
                         uint2 lid [[thread_position_in_threadgroup]])
{
    threadgroup float As[TILE][TILE];
    threadgroup float Bs[TILE][TILE];

    const uint row = gid.y;
    const uint col = gid.x;
    float acc = 0.0f;

    for (uint t = 0; t < (N + TILE - 1) / TILE; ++t) {
        const uint aCol = t * TILE + lid.x;
        const uint bRow = t * TILE + lid.y;
        As[lid.y][lid.x] = (row < N && aCol < N) ? A[row * N + aCol] : 0.0f;
        Bs[lid.y][lid.x] = (bRow < N && col < N) ? B[bRow * N + col] : 0.0f;
        threadgroup_barrier(mem_flags::mem_threadgroup);

        for (uint k = 0; k < TILE; ++k) {
            acc += As[lid.y][k] * Bs[k][lid.x];
        }
        threadgroup_barrier(mem_flags::mem_threadgroup);
    }
    if (row < N && col < N) {
        C[row * N + col] = acc;
    }
}

The Swift side that dispatched it:

let tiles = (n + 15) / 16
enc.setComputePipelineState(pso)
enc.setBuffer(bA, offset: 0, index: 0)
enc.setBuffer(bB, offset: 0, index: 1)
enc.setBuffer(bC, offset: 0, index: 2)
enc.setBytes(&N, length: 4, index: 3)
enc.dispatchThreadgroups(MTLSize(width: tiles, height: tiles, depth: 1),
                         threadsPerThreadgroup: MTLSize(width: 16, height: 16, depth: 1))

The MSL specification defines threadgroup_barrier this way: all threads in a threadgroup executing the kernel "need to execute this function before any thread can continue execution beyond the threadgroup_barrier." The mem_flags::mem_threadgroup argument orders threadgroup memory accesses around the barrier.

Matrix hardware in Metal

MSL exposes cooperative matrix types. The specification says "Metal 2.3 and later support SIMD-group matrix types," with a simdgroup_matrix template in the metal_simdgroup_matrix header, and that operations on these matrices "are executed cooperatively by threads in the SIMD-group," so they must run under uniform control flow. The Feature Set Tables list "SIMD-scoped matrix multiply operations" from the Apple7 family. This is Metal's equivalent of CUDA's WMMA: the next step after shared memory tiling. Apple also ships the Metal Performance Shaders framework, which it describes as "highly optimized compute and graphics shaders," and most users should check it before writing their own kernels.

Metal strengths and costs

Strengths: first class support on every current Apple device, tight integration with Xcode's GPU tools, unified memory on Apple silicon that removes most explicit host to device copies, and run time compilation from source, which kept our correctness test short.

Costs: Apple GPUs only. There is no route from MSL to NVIDIA, AMD or Intel GPUs. Datacentre AI training is not Metal's market, so the AI library ecosystem is smaller than CUDA's or ROCm's.

Vulkan compute: the graphics API's compute half

Vulkan is a cross-vendor graphics and compute API. Khronos announced Vulkan 1.0 on 16 February 2016 as providing "high-efficiency, cross-platform access to graphics and compute on modern GPUs." The current specification is 1.4.366. Compute shaders are usually written in GLSL or HLSL and compiled ahead of time to SPIR-V, the Khronos intermediate representation that the driver consumes.

The tiled kernel in GLSL for Vulkan

// GLSL compute shader for Vulkan: tiled FP32 matmul.
// Follows the GLSL 4.60 and Vulkan specifications; not compiled or run here
// (no glslc or glslangValidator installed on this Mac).
#version 450
layout(local_size_x = 16, local_size_y = 16, local_size_z = 1) in;

layout(std430, set = 0, binding = 0) readonly  buffer BufA { float A[]; };
layout(std430, set = 0, binding = 1) readonly  buffer BufB { float B[]; };
layout(std430, set = 0, binding = 2) writeonly buffer BufC { float C[]; };
layout(push_constant) uniform Params { uint N; } p;

const uint TILE = 16;
shared float As[TILE][TILE];
shared float Bs[TILE][TILE];

void main()
{
    const uint lx  = gl_LocalInvocationID.x;
    const uint ly  = gl_LocalInvocationID.y;
    const uint col = gl_GlobalInvocationID.x;
    const uint row = gl_GlobalInvocationID.y;
    const uint N   = p.N;
    float acc = 0.0;

    for (uint t = 0; t < (N + TILE - 1) / TILE; ++t) {
        const uint aCol = t * TILE + lx;
        const uint bRow = t * TILE + ly;
        As[ly][lx] = (row < N && aCol < N) ? A[row * N + aCol] : 0.0;
        Bs[ly][lx] = (bRow < N && col < N) ? B[bRow * N + col] : 0.0;
        memoryBarrierShared();
        barrier();

        for (uint k = 0; k < TILE; ++k) {
            acc += As[ly][k] * Bs[k][lx];
        }
        barrier();
    }
    if (row < N && col < N) {
        C[row * N + col] = acc;
    }
}

Compile to SPIR-V and dispatch:

# Either compiler works; both produce SPIR-V for vkCreateShaderModule.
glslc -fshader-stage=compute matmul.comp -o matmul.spv
glslangValidator -V matmul.comp -o matmul.spv

// After creating the pipeline, descriptor set and push constants:
vkCmdBindPipeline(cmd, VK_PIPELINE_BIND_POINT_COMPUTE, pipeline);
vkCmdBindDescriptorSets(cmd, VK_PIPELINE_BIND_POINT_COMPUTE, layout, 0, 1, &set, 0, NULL);
vkCmdPushConstants(cmd, layout, VK_SHADER_STAGE_COMPUTE_BIT, 0, sizeof(uint32_t), &N);
vkCmdDispatch(cmd, (N + 15) / 16, (N + 15) / 16, 1);

The GLSL specification says that for compute shaders barrier() may sit inside control flow only if that control flow is uniform, and that making shared writes visible to other invocations "requires the use of both barrier() and a memory barrier." That is why the kernel calls memoryBarrierShared() before the first barrier(). The second barrier protects the tile from being overwritten too early, the same job as CUDA's second __syncthreads().

The limits that make Vulkan different

Vulkan's specification lists minimum guaranteed limits, and they are low because Vulkan runs on phones as well as workstations.

The specification also notes that "the maximum number of invocations that an implementation can support per subgroup is 128." So a portable Vulkan kernel queries VkPhysicalDeviceLimits and VkPhysicalDeviceSubgroupProperties at start up and picks a variant, usually by setting workgroup size through specialisation constants (GLSL's local_size_x_id).

Tensor core style hardware is reachable through the VK_KHR_cooperative_matrix extension, which depends on SPV_KHR_cooperative_matrix and lists a last modified date of 3 May 2023. Support varies by driver and device, so check before you plan around it.

Vulkan compute strengths and costs

Strengths: very wide reach across desktop and mobile GPUs from many vendors, including Android, Windows and Linux devices. It is the natural choice when compute sits alongside rendering in the same application, or when you must ship to consumer hardware from many vendors.

Costs: the most host side code of any model here: instances, devices, queues, memory allocation, descriptor sets, pipeline layouts and command buffers before the first dispatch. The math library ecosystem is thin compared with CUDA or ROCm, and the guaranteed limits are low enough that you must query and branch.

Triton: tiles in Python, compiled for you

Triton began as a research project. Philippe Tillet, H. T. Kung and David Cox presented "Triton: An Intermediate Language and Compiler for Tiled Neural Network Computations" at MAPL on 22 June 2019. The paper describes "a language and compiler centered around the concept of tile, i.e., statically shaped multi-dimensional sub-arrays," and reports matrix multiplication and convolution kernels "on par with hand-tuned vendor libraries (cuBLAS / cuDNN)" in its experiments. Today Triton is a Python embedded language. Its project README lists supported platforms as Linux, with NVIDIA GPUs of compute capability 8.0 and above and AMD GPUs on ROCm 6.2 and above, and CPUs under development.

The tiled kernel in Triton

This is a trimmed version of the official Triton matrix multiplication tutorial, kept in FP32 output to match the other six and without the tutorial's grouped launch order and autotuning.

# Triton: tiled matmul, adapted from the official tutorial.
# Not run for this article (requires Linux with a supported NVIDIA or AMD GPU).
import torch
import triton
import triton.language as tl

@triton.jit
def matmul_kernel(a_ptr, b_ptr, c_ptr, M, N, K,
                  stride_am, stride_ak, stride_bk, stride_bn,
                  stride_cm, stride_cn,
                  BLOCK_M: tl.constexpr, BLOCK_N: tl.constexpr,
                  BLOCK_K: tl.constexpr):
    pid_m = tl.program_id(axis=0)
    pid_n = tl.program_id(axis=1)
    offs_m = pid_m * BLOCK_M + tl.arange(0, BLOCK_M)
    offs_n = pid_n * BLOCK_N + tl.arange(0, BLOCK_N)
    offs_k = tl.arange(0, BLOCK_K)
    a_ptrs = a_ptr + offs_m[:, None] * stride_am + offs_k[None, :] * stride_ak
    b_ptrs = b_ptr + offs_k[:, None] * stride_bk + offs_n[None, :] * stride_bn

    acc = tl.zeros((BLOCK_M, BLOCK_N), dtype=tl.float32)
    for k in range(0, tl.cdiv(K, BLOCK_K)):
        k_left = K - k * BLOCK_K
        a = tl.load(a_ptrs, mask=(offs_m[:, None] < M) & (offs_k[None, :] < k_left), other=0.0)
        b = tl.load(b_ptrs, mask=(offs_k[:, None] < k_left) & (offs_n[None, :] < N), other=0.0)
        acc = tl.dot(a, b, acc)
        a_ptrs += BLOCK_K * stride_ak
        b_ptrs += BLOCK_K * stride_bk

    c_ptrs = c_ptr + offs_m[:, None] * stride_cm + offs_n[None, :] * stride_cn
    tl.store(c_ptrs, acc, mask=(offs_m[:, None] < M) & (offs_n[None, :] < N))

def matmul(a: torch.Tensor, b: torch.Tensor) -> torch.Tensor:
    M, K = a.shape
    _, N = b.shape
    c = torch.empty((M, N), device=a.device, dtype=torch.float32)
    grid = (triton.cdiv(M, 64), triton.cdiv(N, 64))
    matmul_kernel[grid](a, b, c, M, N, K,
                        a.stride(0), a.stride(1), b.stride(0), b.stride(1),
                        c.stride(0), c.stride(1),
                        BLOCK_M=64, BLOCK_N=64, BLOCK_K=32)
    return c

Compared with the CUDA version, there is no thread index, no shared memory declaration and no barrier. Each program instance owns a 64 by 64 tile of C, loads 64 by 32 and 32 by 64 tiles of A and B as whole blocks, and calls tl.dot. The compiler decides how many threads run the program, how the tiles are staged through shared memory, and whether tl.dot maps to tensor core instructions. The bounds masks replace the if (row < N) guards.

A precision detail you must know

tl.dot on two FP32 inputs does not necessarily compute in full FP32. The Triton API reference lists an input_precision argument; for NVIDIA it accepts "tf32", "tf32x3", "ieee", "bf16x3" and "bf16x6" with a default of "tf32", and for AMD it lists "ieee", "bf16x3", "bf16x6" and, on CDNA3 only, "tf32". If your numerical tests compare against a strict FP32 reference, set input_precision="ieee" or expect small differences. The CUDA, HIP, SYCL, OpenCL, Metal and GLSL kernels above all use plain FP32 multiply adds.

Autotuning

The full tutorial wraps the kernel in @triton.autotune with a list of block size configurations and a key of ['M', 'N', 'K'], so Triton benchmarks the configurations the first time it sees a new shape and caches the winner. That is how one Triton source adapts to different GPUs. A 2025 paper by Ringlein, Parnell and Stoica argues that combining just in time compilation with broad autotuning enables portable LLM inference kernels, reporting that their approach explores up to 15 times more kernel parameter configurations and in some cases outperforms vendor optimised implementations by up to 230 percent, while cutting kernel code size by 70 times. Those are the authors' measurements on their kernels and hardware; they show what autotuning can do, not what it will do for you.

Triton strengths and costs

Strengths: very high productivity for AI kernels, fusion of operations (the tutorial applies an activation function while the accumulator is still in FP32), one source that runs on NVIDIA and AMD GPUs, and close integration with PyTorch.

Costs: Linux and two GPU vendors today. Less control: if the compiler makes a poor choice for your shape, your options are block sizes, hints and compiler versions. Debugging generated code needs familiarity with the lower layers. And a Triton kernel's performance is a property of the Triton compiler version as much as of your code, which makes pinning that version part of reproducibility.

Side by side: portability, tooling, libraries and vendor support

"Vendors reached" assumes drivers and implementations are installed and maintained on your fleet. The profiler row matters more than it looks: the 2024 Davis et al. study says "Our ability to identify bottlenecks depended heavily on profiling tools" and called line level stall attribution "a crucial capability missing" from one vendor's profiler at the time.

What the performance portability evidence says

"Portable" means the code runs. "Performance portable" means it runs well everywhere. The second is harder to claim, and the field has a metric for it.

The Pennycook metric, worked

Pennycook, Sewall and Lee proposed in 2016 that performance portability be "the harmonic mean of an application's performance efficiency observed across a set of platforms," and that "If the application fails on any measured platform(s), then we define the performance portability to be 0." Efficiency on each platform is either the fraction of peak (architectural efficiency) or the fraction of the best known implementation (application efficiency).

A worked example, using made up efficiencies to show the arithmetic, not results for any real model:

PP for the first row is 3 / (1/0.9 + 1/0.8 + 1/0.5) = 0.688. The harmonic mean punishes the weak platform far harder than the arithmetic mean does (0.243 against 0.600 in the second row), which is the point: a model that is excellent on two vendors and poor on the third is not performance portable for a buyer who needs all three.

What the studies found

Davis et al., 2024 ("Taking GPU Programming Models to Task for Performance Portability"). The authors compared CUDA, HIP, Kokkos, RAJA, OpenMP, OpenACC and SYCL using five proxy applications on NVIDIA V100, A100 and H100 and AMD MI50 and MI250X systems, including Summit, Frontier and Perlmutter. Their summary: "we find that SYCL most consistently achieves performance portability for our tests, followed closely by RAJA and Kokkos," with a moderate preference for OpenMP over OpenACC among directive models. Compared with earlier studies they found "improved performance portability across models, particularly for SYCL." They also flagged weak spots: poor reduction performance in OpenACC and OpenMP, poor reductions on AMD with RAJA, and "consistent difficulty with the compute-bound and register-intensive miniBUDE for all programming models and systems."

Reguly, 2023 ("Evaluating the performance portability of SYCL across CPUs and GPUs on bandwidth-bound applications"). Using DPC++ and hipSYCL/OpenSYCL on an Intel Data Center GPU Max 1100, an NVIDIA A100 and an AMD MI250X, plus three CPUs, the paper reports that "On GPU architectures SYCL on average even slightly outperforms native approaches, while on CPUs it falls behind." It also cautions that SYCL "does not solve all the challenges of performance portability (e.g. needing different algorithms on different hardware)."

Costanzo et al., 2023 (CUDA versus SYCL for Smith-Waterman protein database search). On NVIDIA, AMD and Intel GPUs, the authors report that CUDA and SYCL achieve similar performance on NVIDIA devices, and that the SYCL version showed "remarkable code portability" to AMD and Intel, with architectural efficiency "superior in 3 of the 4 cases tested."

How to read these results

Each study measured specific applications, compilers, drivers and hardware at one point in time. A SYCL result on bandwidth bound proxy apps says little about a register heavy attention kernel. The defensible conclusions are narrower:

1. Portable models are no longer automatically slower. For several classes of application, SYCL and the C++ abstraction libraries reach performance close to native code on NVIDIA and AMD GPUs.
2. Compute bound, register intensive kernels remain hard everywhere, and that is where vendor libraries and vendor specific tuning still earn their keep.
3. Algorithms may need to change between architectures even when the source language does not.
4. Measure on your own workloads with the Pennycook metric, and report which compiler and driver versions you used.

Step by step: porting the CUDA kernel to the other models

Here is the practical sequence for a team with a CUDA code base that needs a second vendor.

Step 1. Inventory what you depend on. List the CUDA libraries you call (cuBLAS, cuDNN, NCCL, Thrust, CUB), the PTX inline assembly or intrinsics you use, and every place you assume a warp of 32. A quick search finds most of the warp assumptions:

# Find hard coded warp sizes and 32 bit lane masks in a CUDA code base.
grep -rnE "0xffffffff|warpSize|__shfl|__ballot|\b32\b.*lane|lane.*\b32\b" --include="*.cu" --include="*.cuh" src/

Step 2. Choose the target model. For AMD only, HIP is the shortest path. For NVIDIA plus AMD plus Intel from one C++ code base, SYCL. For AI kernels that live next to PyTorch, Triton. For Apple devices, Metal. For consumer graphics applications on many vendors, Vulkan.

Step 3. Run the migration tool and read every diff.

# CUDA to HIP: scan first, then convert in place (keeps .prehip backups).
hipexamine-perl.sh src/
hipify-perl --inplace src/kernels/matmul.cu

# CUDA to SYCL with SYCLomatic: generates SYCL sources in an output folder.
c2s --in-root=src --out-root=sycl_out src/kernels/matmul.cu

The flags above follow each project's documentation at the time of writing; check hipify-perl and c2s help output for your installed version.

Step 4. Fix the warp size and index order. Replace literal 32 with warpSize (HIP) or the queried sub group size (SYCL). In SYCL, make sure the fastest varying index is the last dimension of your nd_range.

Step 5. Swap libraries. cuBLAS calls become hipBLAS or rocBLAS calls on AMD, or oneMKL style calls in SYCL. Check numerical tolerances: different libraries may choose different summation orders and reduced precision paths.

Step 6. Test correctness before speed. Compare against a CPU reference on awkward sizes, not just powers of two. Our Metal and OpenCL checks used N = 100 precisely because 100 is not a multiple of 16, which exercises every bounds check in the kernel.

Step 7. Measure with the vendor profiler and compute PP. Record efficiency on each platform and compute the harmonic mean. If one platform drags the number down, fix that platform before declaring the port done.

Step 8. Pin the toolchain. Record driver, toolkit, compiler and library versions in your build and in your SBOM. A port that works with one ROCm or oneAPI release and fails with the next is a supply chain problem, not just a performance one.

When to choose which

If your team ships CUDA and only buys NVIDIA, a portability rewrite costs money and buys nothing until a second vendor appears. If you expect more than one vendor in the next hardware cycle, the evidence makes SYCL or a C++ abstraction layer worth a pilot, and Triton worth one for AI kernels. Our estimate of GPU cost per million tokens shows why hardware choice moves the economics enough that keeping a second vendor open can matter.

Common mistakes

• Hard coding a warp size of 32. AMD documents 64 wide wavefronts on CDNA. Use warpSize or query the sub group size, and use 64 bit lane masks where the hardware needs them.
• Assuming a 1,024 thread group works everywhere. It does on CUDA and on Apple4 and later Metal GPUs, but Vulkan core guarantees only 128 invocations and Apple2 and Apple3 families allow 512. Query the limit.
• Getting the fastest varying index wrong in SYCL. The last dimension varies fastest. Put the column index there or your loads stop coalescing.
• Dropping the second barrier. Without it, a fast thread can overwrite the shared tile while a slow thread is still reading it. The bug appears only under some timings and on some vendors.
• Missing the memory barrier in GLSL. The GLSL specification says visibility of shared writes needs both barrier() and a memory barrier.
• Placing a barrier in divergent control flow. GLSL requires uniform control flow for barrier(); Metal's SIMD group matrix operations require uniform control flow in the SIMD group. Divergent barriers hang or produce undefined results.
• Comparing Triton FP32 output to a strict FP32 reference. tl.dot defaults to TF32 on NVIDIA for FP32 inputs. Set input_precision="ieee" when you need strict FP32.
• Treating one benchmark as proof of portability. Use the Pennycook harmonic mean across every platform you care about, on your own workloads.
• Testing only power of two sizes. Bounds check bugs hide there. Test sizes such as 100 or 1,000.
• Forgetting the toolchain is a dependency. Drivers, compilers and runtime libraries change behaviour between releases. Pin them and track them.

The software supply chain under every model

Whichever model you choose, the kernel is the smallest part of what you ship. A CUDA application depends on a GPU driver, the CUDA runtime, cuBLAS or cuDNN, often NCCL, and a container base image with all of that inside. A HIP application depends on ROCm components; a SYCL application depends on a SYCL runtime plus the back end vendor stack beneath it; a Triton kernel depends on a specific Triton compiler release, PyTorch and the vendor runtime. Each of those is a versioned software component, and vendors publish security bulletins for drivers, toolkits and container tooling just as they do for any other software. Our post on how a GPU is made covers the hardware side; the software side is where most teams have the least inventory.

Where Safeguard fits

Safeguard does not design, program, compile or tune GPUs, and it does not make GPU code faster. Nothing in this article's kernels or performance discussion is something Safeguard does. The fit is one layer up, in the software you build and ship around those kernels: the Python and C++ dependencies of your AI services, and the container images that bundle framework, runtime and library versions.

Safeguard generates and analyses SBOMs in SPDX and CycloneDX formats, performs software composition analysis across package ecosystems such as npm, PyPI, Maven and Go plus container layers, and matches components against vulnerability data from sources including NVD, GHSA, OSV and CISA KEV, with reachability analysis to help prioritise. In practice that means you can scan the repository and the container image of a model serving service with safeguard scan --dir . or safeguard scan --image <name>, see which projects contain a given component version, and apply policies and gates. Kernel mode GPU drivers installed on a host outside your images, and proprietary binaries without package metadata, may need to be recorded in your SBOM by hand; check what a scan actually reports before relying on it for those layers.

Your GPU toolchain and the open source libraries around it carry the same gap as any other dependency, between a fix being published and that fix running in your production images. Rebuilding a training or serving image pulls those dependencies again, which is why the practices below matter for GPU teams as much as for web teams.

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
• 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
• Intel GPU Architecture: Xe Cores, XMX, Ponte Vecchio and oneAPI SYCL
• 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 JavaScript services that wrap your GPU kernels, and /scan . gives you a starting inventory of a serving repository.

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. Teams that keep GPU build hosts and model code on their own network can run scans where that code already lives.

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 projects in your account contain a component version named in a vendor bulletin, which helps decide which training and serving images to rebuild first.

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.

FAQ

Is HIP the same as CUDA?

Almost, at the source level. HIP mirrors CUDA's runtime API and kernel syntax (__global__, __shared__, threadIdx, __syncthreads), so most kernels port with renaming, which HIPIFY automates. The differences are in the hardware it targets (AMD GPUs), warp size (64 on CDNA, 32 on RDNA), libraries and compiler options.

Is SYCL faster or slower than CUDA?

It depends on the application, compiler and hardware. A 2023 study found SYCL on GPUs slightly ahead of native CUDA and HIP on average for bandwidth bound applications, and a 2024 study found SYCL the most consistently performance portable model across its NVIDIA and AMD tests. Compute bound, register heavy kernels remained hard for every model. Measure your own workload.

Is OpenCL dead?

No. The Khronos specification is still maintained; the unified specification v3.1.2 is dated 17 September 2026 and documents OpenCL 3.1. It still reaches the widest range of devices of any model here; for new portable C++ code, SYCL is the other Khronos option.

Can I run CUDA code on an AMD or Intel GPU?

Not directly. CUDA targets NVIDIA GPUs. The supported routes are to port: HIPIFY converts CUDA to HIP for AMD GPUs, and SYCLomatic's c2s converts CUDA to SYCL, which several implementations compile for NVIDIA, AMD and Intel GPUs. Both tools produce code that still needs review and testing.

What is the difference between Metal and Vulkan compute?

Metal is Apple's own API and shading language for Apple GPUs. Vulkan is a Khronos cross-vendor API whose compute shaders are compiled to SPIR-V. Metal has less host code and Apple specific features such as SIMD group matrix types; Vulkan reaches far more hardware but guarantees lower limits, so portable Vulkan code must query them.

Does Triton work on AMD GPUs?

Yes, according to the Triton README, which lists AMD GPUs on ROCm 6.2 and above alongside NVIDIA GPUs of compute capability 8.0 and above, on Linux. CPU support is listed as under development. Check precision defaults on each vendor: the tl.dot documentation lists different input_precision options for NVIDIA and AMD.

What is a warp, a wavefront, a sub group and a SIMD group?

They are names for the same hardware concept: the bundle of threads that executes one instruction together. CUDA calls it a warp (32 threads). AMD calls it a wavefront in ISA terms and a warp in HIP (64 on CDNA, 32 on RDNA). SYCL and OpenCL say sub group, Vulkan says subgroup, and Metal says SIMD group.

What is the best GPU programming model for beginners?

For learning the concepts, CUDA has the most teaching material, and the ideas transfer directly to HIP, SYCL, OpenCL, Metal and Vulkan. For writing useful AI kernels quickly, Triton has the shortest path from idea to working code. On a Mac, Metal is the practical choice because it runs on the hardware you already have.

Why should a security team care which GPU programming model engineers use?

Because each model brings its own stack of drivers, toolkits, compilers, runtime libraries and container images, each versioned and each subject to vendor security bulletins. The model you choose decides which components land in your SBOM and which vendor advisories you must track.

Quick check

Sources

• NVIDIA CUDA Programming Guide
• CUDA Programming Guide: Programming Model
• CUDA Programming Guide: Writing CUDA Kernels
• CUDA Programming Guide: Writing Tile Kernels
• CUDA Programming Guide: CUDA Platform
• CUDA Programming Guide: Intro to CUDA C++
• CUDA Programming Guide: Compute Capabilities
• NVIDIA CUDA Toolkit Archive
• AMD HIP documentation: What is HIP?
• AMD HIP documentation: Programming model
• AMD HIP documentation: Porting CUDA code to HIP
• AMD HIPIFY documentation
• Khronos SYCL 2020 Specification (revision 12)
• Khronos press release: SYCL 2020 Specification (9 February 2021)
• Intel LLVM repository: oneAPI DPC++ compiler
• AdaptiveCpp (formerly hipSYCL and Open SYCL)
• SYCLomatic
• Khronos OpenCL Specification (unified, v3.1.2)
• Khronos press release: OpenCL 1.0 Specification
• Khronos press release: OpenCL 3.0 provisional specifications
• Apple Metal Shading Language Specification, version 4.1
• Apple Metal Feature Set Tables
• Apple Metal developer page
• Khronos Vulkan Specification (1.4.366)
• Khronos press release: Vulkan 1.0 (16 February 2016)
• Khronos VK_KHR_cooperative_matrix extension
• The OpenGL Shading Language, version 4.60
• Tillet, Kung and Cox, Triton: An Intermediate Language and Compiler for Tiled Neural Network Computations (MAPL 2019)
• Triton tutorial: Matrix Multiplication
• Triton API reference: triton.language.dot
• Triton GitHub repository and README
• Pennycook, Sewall and Lee, A Metric for Performance Portability (2016)
• Davis et al., Taking GPU Programming Models to Task for Performance Portability (2024)
• Reguly, Evaluating the performance portability of SYCL across CPUs and GPUs on bandwidth-bound applications (2023)
• Costanzo et al., Comparing Performance and Portability between CUDA and SYCL for Protein Database Search (2023)
• Tsai et al., Preparing Ginkgo for AMD GPUs: A Testimonial on Porting CUDA Code to HIP (2020)
• Karimi, Dickson and Hamze, A Performance Comparison of CUDA and OpenCL (2010)
• Ringlein, Parnell and Stoica, GPU Performance Portability needs Autotuning (2025)

Last reviewed: 25 October 2026. Toolkit versions, supported hardware lists and specification revisions change often; check the linked vendor and Khronos pages before relying on a specific limit or version.