What Is CUDA, and Why Does GPU Programming Create Vendor Lock-In?

Most people meet CUDA for the first time and assume it's a programming language. It isn't, quite. CUDA is NVIDIA's platform for running general-purpose computation on NVIDIA GPUs. It lets you take selected parts of a program and run them on the GPU instead of the CPU.

The easiest way to see what that actually means is to look at a program.

A first CUDA program

#include <cstdio>
#include <cuda_runtime.h>

__global__ void helloFromGPU()
{
    printf("Hello World from the GPU!\n");
}

int main()
{
    printf("Hello World from the CPU!\n");

    helloFromGPU<<<1, 1>>>();

    cudaDeviceSynchronize();

    return 0;
}

Most of this looks like ordinary C++, because it mostly is. CUDA programming is normally done in CUDA C++, which is regular C++ plus a handful of language extensions plus the CUDA runtime libraries.

What makes this file unusual is that it contains code for two different processors. The main() function runs on the CPU. The function marked __global__ runs on the GPU. And this line, which no standard C++ compiler will accept, is what connects them:

helloFromGPU<<<1, 1>>>();

Those triple angle brackets are a CUDA extension. NVIDIA's compiler understands them. g++ will just give you a syntax error.

What __global__ actually does

The __global__ keyword tells the CUDA compiler that this function is a kernel, meaning a function meant to run on the GPU. Marking it doesn't run it. You still have to launch it, and the launch is where you decide how much of the GPU to use.

The two numbers inside the angle brackets are the number of blocks and the number of threads per block. So <<<1, 1>>> gives you one block containing one thread, which is a single GPU thread running the kernel once. Not exactly a workout for a card with thousands of cores, but it's enough to prove the plumbing works.

One detail worth knowing early: a __global__ function must return void. If you need a result back, you write it into GPU memory and copy it out afterwards.

Threads and blocks

If you've worked with operating system processes and CPU threads, CUDA's hierarchy will look familiar but not identical. A kernel launch creates a grid. The grid holds blocks, and each block holds threads.

A block is just a group of GPU threads. Launch helloFromGPU<<<3, 4>>>() and you get three blocks of four threads each, so twelve GPU threads in total, arranged like this:

Grid
├── Block 0 → Thread 0, 1, 2, 3
├── Block 1 → Thread 0, 1, 2, 3
└── Block 2 → Thread 0, 1, 2, 3

Notice that thread numbering restarts inside every block. Thread 2 of block 0 and thread 2 of block 1 are different threads with the same local number, which is a problem the moment you want each thread to handle a different piece of data.

CUDA gives you three built-in variables to sort this out: threadIdx.x for the thread's position in its block, blockIdx.x for the block's position in the grid, and blockDim.x for how many threads a block holds. Put them together and you get a unique ID for every thread in the launch:

int id = blockIdx.x * blockDim.x + threadIdx.x;

This is probably the single most important line in beginner CUDA. For the <<<3, 4>>> launch above, it produces IDs 0 through 3 in the first block, 4 through 7 in the second, and 8 through 11 in the third. Every thread now knows exactly which element of the data it owns.

Why blocks exist at all

Blocks aren't just a tidy way to draw diagrams. Threads inside the same block get real privileges. They can share fast on-chip memory, they can synchronise with each other, and they can divide up one chunk of a larger problem between them. Threads in different blocks generally can't do any of that, and you should assume blocks run independently and in no particular order. Newer NVIDIA architectures do add a middle layer called thread block clusters, which lets a small group of blocks cooperate.

What happens if you ask for too many threads

Try this:

helloFromGPU<<<10000, 10000>>>();

The block count is fine. The thread count isn't. Current NVIDIA GPUs cap a block at 1,024 threads, and asking for 10,000 will simply fail.

The nasty part is that it fails quietly. A kernel launch doesn't throw, and if you don't check for errors you'll sit there wondering why nothing printed. So check:

helloFromGPU<<<10000, 10000>>>();

cudaError_t error = cudaGetLastError();

if (error != cudaSuccess) {
    printf("Launch error: %s\n", cudaGetErrorString(error));
}

cudaDeviceSynchronize();

Now you get told what went wrong:

Launch error: invalid configuration argument

Change it to <<<10000, 1024>>> and the launch is legal. That's 10,240,000 threads, which sounds absurd until you realise the GPU never intends to run them all at once.

Real kernels do arithmetic. Here's the shape of an honest one:

__global__ void addArrays(
    const int* a,
    const int* b,
    int* result,
    int n
)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;

    if (i < n) {
        result[i] = a[i] + b[i];
    }
}

Each thread adds exactly one pair of numbers. The if (i < n) guard matters more than it looks. Your array size will almost never divide evenly into your block size, so you'll usually launch slightly more threads than you have data. Without that check, the extra threads write past the end of the array, and on a GPU that means silent memory corruption rather than a helpful crash.

That's the pattern for basically all of this: take a big problem, split it into many small identical tasks, and hand one task to each thread.

Is there a process running inside the GPU?

Not in the operating system sense. Your CUDA program is a normal Linux process. The OS starts it on the CPU, and that CPU process is what pushes work onto the GPU. Nothing is "running" on the GPU independently of it.

What the NVIDIA driver does create for your process is a CUDA context. The context holds the GPU-side state that belongs to your program: allocated GPU memory, loaded kernel code, execution state, streams and events, and device configuration. When your process exits, the context goes with it.

If you want a rough mapping between the two worlds:

Linux process  → CUDA context
CPU thread     → launches GPU work
CUDA block     → group of GPU threads
CUDA thread    → one instance of the kernel running

Rough is the operative word. Don't push the analogy hard.

So is CUDA a language or not?

CUDA is a platform and a programming model. The language people write in is CUDA C++, and that's only one piece of a much larger box:

• The programming model: grids, blocks, threads, kernels.
• C++ language extensions: __global__, __device__, and the <<<blocks, threads>>> launch syntax.
• Compiler tooling: nvcc, the compiler driver.
• Runtime APIs: cudaMalloc, cudaMemcpy, cudaFree, cudaDeviceSynchronize, and the rest.
• Accelerated libraries: cuBLAS, cuDNN, cuFFT, NCCL, TensorRT.

That mix is why "CUDA is just C++ with a library" doesn't hold up. A library function like cudaDeviceSynchronize() really is an ordinary function call. But __global__ and kernel<<<1, 1>>>() are compiler-level extensions, and no amount of #include will teach g++ what they mean. That's why CUDA source files use the .cu extension and get built with:

nvcc program.cu -o program

It's also worth saying that compilation happens on the CPU. CUDA doesn't compile C++ on the GPU. What nvcc does is split your file, send the host code to a normal C++ compiler, compile the device code for the GPU, and bundle both into one executable. At runtime the CPU starts up, prepares data, launches the kernel, and later collects the results.

The CUDA Toolkit and the Runtime API

The CUDA Toolkit is the development kit you install to build any of this. It ships the nvcc compiler driver, the CUDA headers, the Runtime API, the maths and AI libraries, plus debuggers, profilers and sample code.

The Runtime API is the part your CPU code actually calls to boss the GPU around. The common ones are cudaMalloc, cudaFree, cudaMemcpy, cudaDeviceSynchronize, cudaGetLastError and cudaGetDeviceProperties. Between them you can allocate GPU memory, move data across, wait for work to finish, ask what hardware you're on, and catch errors.

A typical program is a sandwich:

int main()
{
    cudaMalloc(...);      // reserve GPU memory
    cudaMemcpy(...);      // copy input up to the GPU

    kernel<<<blocks, threads>>>();

    cudaDeviceSynchronize();  // wait for it to finish

    cudaMemcpy(...);      // copy results back down
    cudaFree(...);        // release GPU memory
}

All of that runs on the CPU. It reaches the hardware through several layers:

Your application
        ↓
CUDA Runtime API      (cudaMalloc, cudaMemcpy, ...)
        ↓
CUDA Driver API       (cuMemAlloc, cuLaunchKernel, ...)
        ↓
NVIDIA kernel-mode driver
        ↓
NVIDIA GPU

Two things often get conflated there. The CUDA Driver API is a lower-level programming interface that you can call directly if you want fine control. The kernel-mode driver is the actual operating system driver talking to the hardware. Most people only ever touch the Runtime API, which is the friendlier of the two.

As for what the toolkit itself is written in, there's no single answer, because it isn't a single program. The headers and public interfaces are C and C++. nvcc is built on top of LLVM, which NVIDIA has been fairly open about. Plenty of the rest is closed source, so anything more specific is guesswork.

While we're here, "native software" doesn't mean somebody wrote it in assembly. It means the source was compiled into machine instructions that run directly on the target CPU and OS, rather than running inside an interpreter or a virtual machine.

What companies actually use CUDA for

The short version: anything with a mountain of computation that breaks cleanly into many similar operations.

Right now the dominant use is AI. The stack usually looks like Python at the top, PyTorch or TensorFlow underneath, CUDA libraries under that, and an NVIDIA GPU at the bottom. Most AI engineers never write a kernel. They write:

model = model.to("cuda")

and PyTorch handles everything below that line.

Beyond AI, the list is long. Video encoding, decoding, filters and AI upscaling. Weather modelling, fluid and molecular simulation, medical imaging, computational chemistry. Robotics work like camera processing, object tracking, mapping and sensor fusion. In finance, risk simulation and option pricing. In data systems, vector similarity search, big aggregations, sorting and graph processing.

Different fields, same shape underneath. Thousands of independent calculations, thousands of GPU threads, results in parallel.

How deep companies go

Not everyone uses CUDA the same way, and the depth matters a lot for what comes next.

At the top are the framework users, working in PyTorch, TensorFlow, JAX or off-the-shelf rendering and database tools. Many of them never write a line of CUDA.

In the middle are teams calling CUDA libraries directly: cuBLAS for linear algebra, cuDNN for deep learning primitives, TensorRT for optimised inference, NCCL for multi-GPU communication, cuFFT for Fourier transforms.

At the bottom are the specialists writing .cu files by hand, tuning memory access patterns, thread organisation, shared memory usage, occupancy and multi-GPU communication, sometimes reaching for hardware-specific instructions.

Here's the part that matters. The further down that list a company goes, the harder it becomes to leave.

Moving from NVIDIA to AMD or Intel

AMD

NVIDIA GPUs run CUDA. AMD GPUs run ROCm. The pieces line up reasonably well:

NVIDIA          AMD
CUDA             ROCm
CUDA C++         HIP C++
nvcc             hipcc
cuBLAS           rocBLAS
cuDNN            MIOpen
NCCL             RCCL

A compiled CUDA binary won't run on an AMD card. What the migration costs you depends entirely on how deep your code sits.

If your application is mostly PyTorch, you may get away with very little. Install the ROCm build of PyTorch instead of the CUDA build and the Python code often runs unchanged. You still have to test it properly, because "it runs" and "it runs as fast" are different claims, but the code changes can be small.

If you call the CUDA API directly, you'll be doing translation work. cudaMalloc becomes hipMalloc, cudaMemcpy becomes hipMemcpy, and so on. AMD ships a tool called HIPIFY that automates a good chunk of this, since HIP was deliberately designed to look like CUDA.

If you have custom kernels, expect real effort. The syntax is similar enough to lull you, but the hard parts are elsewhere: third-party libraries that only exist for CUDA, optimisations tuned for NVIDIA's memory hierarchy, anything built on Tensor Cores, pre-compiled CUDA extensions in your dependency tree, and differences in how the two architectures schedule work.

The pattern most teams report is that converting the code is manageable. Getting the performance back is the expensive part.

Intel

Intel's accelerator story is oneAPI, and the language is SYCL:

AMD              Intel
ROCm             oneAPI
HIP C++          SYCL C++
hipcc            Intel's SYCL compiler
rocBLAS          oneMKL
MIOpen           oneDNN
RCCL             oneCCL

This jump feels bigger than the last one, because HIP imitates CUDA's style while SYCL uses standard C++ abstractions instead. Work gets submitted to a queue and the kernel body is a lambda:

queue.parallel_for(
    sycl::range<1>(totalThreads),
    [=](sycl::id<1> index) {
        int i = index[0];
        result[i] = a[i] + b[i];
    }
);

Look past the syntax and it's the same idea you started with. Create a lot of parallel workers, give each one an index, have each one handle its own slice of data. The concept transfers cleanly. The toolchain, the libraries and the tuning knowledge don't.

Cross-vendor options, and their limits

So every major vendor has its own stack: CUDA for NVIDIA, ROCm and HIP for AMD, oneAPI and SYCL for Intel. There are genuine cross-platform options, though. PyTorch and TensorFlow both abstract over backends. SYCL, OpenCL, OpenMP GPU offloading and Vulkan Compute all aim at portability from different angles.

PyTorch is the clearest example, since the same model code can target a CUDA backend on NVIDIA, a ROCm backend on AMD, or an XPU backend on Intel. For a lot of teams that's enough, and it's why the AI world isn't as locked in at the application layer as people sometimes assume.

But an abstraction layer can only hide so much. Underneath, the implementations are still completely different, and the further you get from the common path the more that shows.

Is the lock-in deliberate?

Partly technical, partly commercial. Both halves are real.

The technical half. GPU architectures genuinely differ, and not in trivial ways. They vary in instruction sets, memory systems, scheduling, the size of a thread group, matrix acceleration hardware, interconnects, compiler behaviour and supported data types. NVIDIA has warps and Tensor Cores. AMD has wavefronts and Matrix Cores. Intel has subgroups and XMX engines. These aren't different names for the same thing.

A common platform can cover the basics without much trouble. Allocate memory, copy memory, launch parallel work, synchronise. The specialised hardware is where it breaks down, and a universal platform then has three options: ignore vendor-specific features, expose a slower lowest-common-denominator abstraction, or add vendor-specific extensions. The moment it picks the third option, some of the lock-in walks straight back in.

That's the trade-off nobody has solved. More portability means less hardware-specific optimisation. More optimisation means more vendor-specific code.

The commercial half. Vendors also do very well out of strong software ecosystems, and NVIDIA has run this loop better than anyone. Good tools and libraries attract developers. Developers build applications on CUDA. Those applications become expensive to move. Customers keep buying NVIDIA hardware. Repeat.

NVIDIA has little reason to make CUDA run beautifully on competitors' cards. Doing so would help those competitors sell hardware, erode NVIDIA's main differentiator, and commit NVIDIA to supporting silicon it doesn't control. AMD and Intel build ecosystems around their own hardware too, though both lean harder on portability and open standards, which is at least partly a position you take when you're not the incumbent.

Wrapping up

CUDA isn't a language or a library. It's NVIDIA's full GPU computing platform, and a CUDA C++ application pulls in CPU-side C++, GPU kernels, compiler extensions, runtime APIs and NVIDIA's libraries all at once. The CPU process launches kernels, and those kernels are organised as a grid of blocks of threads.

Companies use it for AI, video, simulation, robotics, finance and data work. It targets NVIDIA hardware. AMD and Intel have their own platforms, and cross-platform frameworks soften the problem without eliminating it, because you cannot fully hide architectural differences behind an abstraction.

A universal GPU programming model is technically possible. Identical features and identical performance across every GPU is a much harder promise, and vendors have every reason to keep strengthening their own platforms, because software ecosystems are how they sell hardware.

The deeper your application reaches into one vendor's libraries and optimisations, the more expensive it becomes to leave. That's worth knowing before you write the code, not after.