Is CUDA a programming language? Strictly speaking, no. NVIDIA describes CUDA as a general-purpose parallel computing platform and programming model for its GPUs. What developers actually write, CUDA C++, is a small set of extensions to C++ plus a runtime library, and NVIDIA’s nvcc compiler driver compiles it through an intermediate representation called PTX. The answer matters for more than terminology. The language layer is C++ with a minimal set of extensions. The platform layer (proprietary compilers, runtime and licence) is the part tied to NVIDIA.

If you came here for the acronym, the page on what CUDA stands for covers the name and where it came from. This article covers what sits underneath it.

Is CUDA a programming language, or C++ with extensions?

A hierarchy of thread groups, shared memories and barrier synchronisation form the core abstractions, and NVIDIA’s programming guide says they are exposed to the programmer as a minimal set of language extensions. The nvcc documentation describes GPU code as written in “a language that is essentially C++”, with annotations that separate device functions from host code and mark the different kinds of GPU memory. Both describe C++ with extensions, not a separate language.

C++ standard support follows the host compiler’s dialect. Passing nvcc -std=c++<version> turns on the features of that C++ standard and invokes the host compiler with the same dialect, subject to documented restrictions. CUDA also ships libcu++, an implementation of the C++ Standard Library that works on both host and device, and adds CUDA-specific extras such as synchronisation primitives and C++ wrappers for PTX instructions.

The restrictions show where CUDA C++ stops being plain C++:

• Run-time type information and exceptions are not supported in device code, so typeid and dynamic_cast are unavailable there.
• Kernel parameters are passed through constant memory, and their total size is limited to 32,764 bytes.
• Non-trivially copyable kernel arguments launched from host code are allowed, but they are not processed according to the standard C++ model.
• The address of a __device__ function cannot be taken in host code.

There is also a case for calling CUDA-C two languages that share one syntax. A formal semantics of CUDA-C (arXiv, 2012) built on an existing C semantics by adding roughly 400 rewrite rules to the original 1,163. It used hooks so that the semantics knows whether it is giving meaning to a term in the CPU language or in the GPU language. The same authors set this apart from PTX, which is used only to code GPU kernels and does not cover CPU-GPU interactions such as memory allocation, transfers and streams.

The layers people mean when they say “CUDA”

Layer	What it is	Evidence
Programming model	Host and device, kernels, the thread hierarchy; NVIDIA defines it independently of any language	NVIDIA programming guide
Language extensions	Execution space specifiers such as __global__ and __device__, plus the <<<...>>> launch syntax	CUDA C++ extensions appendix
Compiler driver	nvcc, which hides the splitting, compiling and merging steps of CUDA compilation	nvcc documentation
Intermediate representation	PTX, assembly for a virtual GPU architecture	nvcc documentation
Runtime and driver APIs	The runtime API, built on the lower-level CUDA driver API	NVIDIA programming guide
Libraries	cuBLAS, cuFFT, cuSPARSE and NCCL, also reachable from Fortran	CUDA Fortran guide
Licence	NVIDIA’s SDK licence agreement with a CUDA Toolkit supplement	CUDA EULA

From a .cu file to GPU code: what nvcc actually does

A CUDA source file mixes conventional C++ host code with GPU device functions. nvcc separates the two, compiles the device functions with NVIDIA’s proprietary compilers and assembler, compiles the host code with an available C++ host compiler, and embeds the compiled GPU code as fatbinary images in the host object file. The device pass produces CUDA binary (cubin), PTX or cuda_tile intermediate code, all placed in a fatbinary. A second host pass then embeds that fatbinary and turns the CUDA-specific C++ extensions into standard C++ constructs. That second pass is the most concrete evidence for the “extension, not language” answer: what the host compiler finally sees is ordinary C++.

You can write kernels directly in PTX, but NVIDIA notes that a high-level language such as C++ is usually more effective, and either way nvcc must compile kernels to binary code before they run on the device.

The two-architecture model behind PTX

nvcc uses a two-stage compilation model so that applications stay compatible with future GPU generations. The first stage targets PTX, which nvcc’s documentation calls assembly for a virtual GPU architecture: a GPU defined entirely by the features it offers, with a largely generic instruction set, always stored as text. For that reason every nvcc command involves two architectures, a virtual one for the intermediate code and a real one naming the processor the code will run on.

If you specify only a virtual architecture, for example --gpu-architecture=compute_90 --gpu-code=compute_90, nvcc postpones assembly until run time, when the actual GPU is known. A fatbinary can hold several translations of the same source, and the CUDA driver picks the most appropriate one at launch. When the runtime creates its context for a device, it just-in-time compiles device code if needed and loads it into device memory, without any action from the developer.

This choice affects whether a binary keeps working on future GPU generations. Shipping PTX lets assembly wait until run time, when the target GPU is known. Shipping only real-architecture code ties the build to the architectures you listed.

nvcc also supports architecture-specific code paths. The __CUDA_ARCH__ macro holds the virtual architecture being compiled for during each stage-one compilation, and NVIDIA documents it as the way to specialise code for particular GPU architectures, for performance or to reach architecture-specific instructions. Every one of those branches is a place where code becomes NVIDIA-specific, and often generation-specific too.

nvcc arrives as part of a versioned SDK, which is a separate thing from both the programming model and the GPU driver; CUDA vs CUDA Toolkit explains why that distinction matters when you pin versions.

What is a CUDA kernel, and how does it reach the GPU’s threads?

A CUDA kernel is a C++ function that, when called, is executed N times in parallel by N CUDA threads, rather than once like a regular C++ function. You mark it with the __global__ specifier, which tells the compiler to build the function for the GPU so it can be invoked by a kernel launch. Applications always start on the CPU. Host code uses CUDA APIs to copy data, start GPU work and wait for it to finish.

Kernels are launched with triple chevron notation, a CUDA C++ language extension. Each kernel call carries an execution configuration in the form <<<grid_dim, block_dim, dynamic_smem_bytes, stream>>>, and the product grid_dim.x * grid_dim.y * grid_dim.z equals the number of thread blocks launched. Inside the kernel, each thread works out which piece of data it owns from its block and thread indices.

Worked example: one million elements, 256 threads per block

__global__ void scale(float* x, float a, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) x[i] *= a;
}

// host code
int threads = 256;
int blocks  = (n + threads - 1) / threads;
scale<<<blocks, threads>>>(d_x, 2.0f, n);

Then blocks = ⌈1,000,000 / 256⌉ = ⌈3,906.25⌉ = 3,907, which launches 3,907 × 256 = 1,000,192 threads. The last 192 threads have no element to process, and that is why the if (i < n) guard is there.

The grid, thread blocks and streaming multiprocessors (SMs) are the CUDA GPU architecture that both of CUDA’s kernel styles build on. Each thread block runs on a single SM, where its threads share that SM’s on-chip shared memory and can synchronise with one another. A few rules come with that mapping. __syncthreads() is allowed inside conditional code only if the condition evaluates the same way across the whole block; otherwise the kernel may hang. CUDA also assumes a weakly ordered memory model, so unsynchronised reads and writes to the same location are undefined behaviour.

Kernels are not the only GPU functions. __device__ marks a GPU function that other device or kernel functions can call, and __host__ __device__ tells the compiler to generate both CPU and GPU code for the same function. Memory specifiers such as __shared__, __constant__ and __managed__ control where a variable lives. More recent versions add tile programming alongside the classic SIMT style, chosen per kernel and running on the same device memory, and from CUDA 13.3 both kinds of code can sit in the same translation unit. For the wider picture of why so many threads are worth it, GPU parallel computing explained covers the execution model beyond CUDA.

Which languages can target CUDA besides C++, and what do you give up?

NVIDIA presents the CUDA programming model separately from any language and says its concepts apply in every supported one. Alongside C++, the programming guide lists Fortran, DirectCompute and OpenACC as supported routes. Each route sits on the same platform, so each inherits the same NVIDIA dependency unless it also has a separate back end for other vendors.

Route	How it reaches the GPU	Beyond NVIDIA?	What you give up compared with CUDA C++
CUDA Fortran	A small set of Fortran extensions; kernels are attributes(global) subroutines called with chevron syntax	No: supported by the NVIDIA Fortran compiler	Dependence on the NVIDIA Fortran compiler; one exascale team moved its GPU Fortran to C++, citing weak Fortran compiler support on GPUs as one of several reasons
CUDA Python	NVIDIA’s tools and libraries for reaching the CUDA platform from Python	No: built on NVIDIA’s CUDA platform	NVIDIA says Python and C++ are equal first-class citizens as of CUDA 13.3, with feature parity as its aim going forward
Numba (numba.cuda)	Compiles a restricted subset of Python into CUDA kernels	Disputed (see below)	Only part of Python is supported, though thread indexing is simpler through cuda.grid
PyCUDA	Exposes the CUDA API to Python	No: NVIDIA GPUs only	Any route to AMD or Intel GPUs
Julia (CUDA.jl)	@cuda compiles kernels on first use, through LLVM’s NVPTX back end to PTX	Yes, through AMDGPU.jl, oneAPI.jl and Metal.jl	A dependency on Julia’s code generation rather than nvcc’s
Kokkos and similar layers	C++ templates over CUDA or HIP back ends	Yes	Templates fix choices such as threads per block and transfer overlap early, which hinders kernel-specific tuning

Fortran users have a middle path. Fortran 2003’s C interoperability lets code call both CUDA and HIP without relying on non-standard extensions. For directive-based offload, CUDA vs OpenACC compares directives with hand-written kernels.

The sources disagree on Numba and AMD. A 2023 study of Julia, Python/Numba and Kokkos on exascale nodes reports Python/Numba support for AMD GPUs as deprecated and lists it as unsupported in its test setup. The ENCCS GPU-programming training material says Numba supports both NVIDIA and AMD. We give more weight to the study, because it records the versions it tested.

Julia’s published benchmarks show a high-level language performing close to CUDA C. Besard et al. (2017) measured an average slowdown of 0.50% against CUDA C kernels compiled with nvcc on Rodinia benchmarks, and found the generated PTX identical to the equivalent CUDA C apart from inlining decisions. The 2023 exascale study reports a mixed result. It found Julia’s AMDGPU.jl comparable to HIP on AMD, but reported a remaining gap on NVIDIA A100 GPUs for single-precision cases that needs further investigation into the PTX being generated.

TechnoLynx’s view is that what CUDA C++ offers over these routes is access to lower-level features rather than built-in speed. On that view, vendor-specific APIs are not inherently faster: the same code performs essentially the same, and the benefit of the lower-level features they expose only arrives with real commitment, which is not always worth making. Those features are concrete things: __CUDA_ARCH__-specialised paths, or DPX instructions whose hardware acceleration depends on compute capability. Whether a kernel needs them is something profiling tells you, not a rule of thumb.

Is CUDA free, and what does being tied to NVIDIA cost?

The CUDA Toolkit’s use is governed by NVIDIA’s SDK licence agreement plus a CUDA Toolkit supplement. The licence lists the toolkit files you may ship with your own applications, such as the CUDA runtime library, and asks organisations whose distribution needs fall outside those terms to contact NVIDIA. The costs the sources document are the constraints that come with the platform.

The first constraint is hardware. An LLNL paper on breaking vendor lock-in notes that, as a proprietary programming language, CUDA can traditionally be compiled only to NVIDIA GPUs, and that programming accelerators in vendor-specific languages such as CUDA is not portable and leads to vendor lock-in.

The second constraint is legal. Tom’s Hardware reported in March 2024 that the EULA installed with CUDA 11.6 and later includes a clause forbidding reverse engineering, decompiling or disassembling output generated with SDK elements in order to translate it to a non-NVIDIA platform. According to the same report, the restriction had appeared in NVIDIA’s online licence terms since 2021. The same report notes that recompiling CUDA programs for other hardware remains legal. In engineering terms, the permitted way off NVIDIA starts from source code, not from compiled binaries.

HIP and other exits from CUDA

HIP is AMD’s C++ runtime API and kernel language, and its runtime API closely mirrors CUDA’s. HIPIFY automatically translates CUDA code to HIP, most CUDA libraries have a corresponding HIP library, and HIP keeps the triple-chevron launch syntax. HIP also replaces architecture-number checks on __CUDA_ARCH__ with feature-specific macros. Exascale teams describe both styles of port: converting all CUDA code with hipify and adding adapters to libraries such as rocBLAS, or keeping a thin abstraction layer that calls either cudaSetDevice() or hipSetDevice() depending on build configuration. HIP can target both AMD and NVIDIA, but the code has to be ported or kept compliant with hipify’s requirements. Research compilers go further. LLNL’s prototype compiles unmodified CUDA source for both NVIDIA and AMD GPUs through an OpenMP offload runtime, though the paper presents it as a prototype.

How much an exit costs depends on what the code actually uses, more than on the fact that it is CUDA:

What the code depends on	What leaving involves
Frameworks only, no custom kernels	Mostly a framework and build question; the CUDA dependency sits inside the framework’s back end
Custom kernels using the core runtime API	A source translation of the kind HIPIFY automates
NVIDIA libraries such as cuBLAS or cuFFT	Mapping each call to its HIP counterpart, which exists for most CUDA libraries
__CUDA_ARCH__ paths, inline PTX, architecture-specific instructions	Rewriting and re-tuning those paths for each target, which is the expensive part
Compiled binaries only, no source	The EULA clause Tom’s Hardware reported bars translating compiled output to target a non-NVIDIA platform

How CUDA compares with portable APIs when you start fresh is covered in CUDA vs OpenCL vs SYCL, and CUDA vs AMD looks at the HIP route specifically. The useful question in all of these is how wide the CUDA commitment really is, compared with OpenCL, SYCL or HIP for your workload. Porting CUDA kernels to HIP, or keeping one codebase fast on more than one vendor, is work TechnoLynx’s GPU performance engineering team takes on.

Measuring how much of a codebase belongs to NVIDIA

A .cu file mixes conventional C++ host code with device functions. Search the code for __CUDA_ARCH__ branches, inline PTX, driver-API calls and architecture-specific intrinsics. Those lines are where the code depends on NVIDIA’s platform rather than on C++. If there are few of them, only a small share of the code is NVIDIA-specific. If they carry the hot paths, whether those features pay is a question profiling answers. You have to settle that before arguing about whether to stay or port.