NVIDIA CUDA on Ubuntu 26.04

NVIDIA CUDA on Ubuntu 26.04

based on An Even Easier Introduction to CUDA

If you have an NVIDIA GPU you can use CUDA, a parallel computing platform for writing high-performance applications using thousands of parallel threads on GPUs. Everything on this article also works on Ubuntu flavors like Kubuntu.

Prerequisites

The first step is to install the proprietary NVIDIA driver. See which one is the recommended one with:

You can automatically install it with:

Optionally, you can verify which driver packages are installed and see that the recommended one was installed:

You'll also need the CUDA Toolkit. On its website you'd find versions for older Ubuntu releases but now on 26.04 you can install it directly with:

For profiling our application, we'll need NVIDIA Nsight Systems. Download and install the package for your system which is probably the Linux x8664 .deb installer.

Now download this wrapper script for the nsys command that makes profiling your CUDA app easier. Put it somewhere on your $PATH or in the directory where you'll be working.

Key Concepts A kernel in CUDA is a function that the GPU can run. You specify it with global. Unified Memory in CUDA provides a single memory space accessible by all CPUs and GPUs in your system. This memory allocation returns a pointer that can be accessed by host (CPU) code or device (GPU) code. If we know what memory is needed by the kernel we can prefetch it to make sure the data is on the GPU before the kernel needs it. This avoids a situation where you have multiple page faults and the hardware migrates the pages to the GPU memory when the faults occur. CUDA GPUs have many parallel processors grouped into Streaming Multiprocessors (SMs). Each SM can run multiple concurrent thread blocks, but each block runs on a single SM. Together, the blocks of parallel threads make up what is known as the grid. When launching a kernel you specify the number of thread blocks and block size with >> notation meaning: $$ A $$ thread blocks $$ B $$ threads in a block (in multiples of $$ 32 $$) CUDA C++ provides keywords that let kernels get the indices of the running threads. gridDim.x is the number of blocks in the grid. blockIdx.x is the index of the current thread block in the grid. blockDim.x is the number of threads in the block. threadIdx.x is the index of the current thread within its block. So in order to process an array of $$ 2^{20} $$ elements with a block size of $$ 256 $$ threads we have $$ 2^{20}/256 = 4096 $$ blocks.

Figure: Grid, Block and Thread indexing in CUDA kernels (one-dimensional) from the NVIDIA blog

Example Let's use parallel processing on the GPU to add the elements of two arrays. Save this program in a file called add_grid.cu

Compile it:

Profile it:

That's it! On my NVIDIA RTX 5070 Ti the add operation for ~1M elements takes only 21 microseconds (µs) thanks to the parallelism on the GPU and the memory prefetching. This is compared to ~85,000 microseconds (µs) in a version of this code that uses only 1 thread on the GPU (without parallelism) and no memory prefetching. A huge performance boost!

What results do you get on your machine?