CUDA

Guidance last updated: 22 October 2026

CUDA (Compute Unified Device Architecture) is a parallel computing platform and application programming interface (API) model created by NVIDIA. It allows software developers to use a CUDA-enabled graphics processing unit (GPU) for general purpose processing, an approach known as General Purpose GPU (GPGPU) computing.

Usage

You need to first request one or more GPUs within an interactive session or batch job on a worker node.

To request one A100 GPU for a batch job you would include all the following in the header of your submission script:

#SBATCH --partition=gpu
#SBATCH --nodes=1
#SBATCH --gres=gpu:a100:1

Note

See Using GPUs on Stanage for more information on how to request a GPU-enabled node for an interactive session or job submission.

You then need to ensure a version of the CUDA library (and compiler) is loaded. As with much software installed on the cluster, versions of CUDA are activated via the ‘module load’ command.

To load one of the currently available CUDA versions you can run one of the following commands:

module load CUDA/13.2.0
module load CUDA/12.9.1
module load CUDA/12.8.0
module load CUDA/12.4.0
module load CUDA/12.1.1

module load CUDA/13.2.0
module load CUDA/12.9.1
module load CUDA/12.8.0
module load CUDA/12.4.0

CUDA requires a supported C++ host compiler. For the examples below, load:

module load CUDA/13.2.0 GCC/15.2.0

Check the toolkit and driver separately:

nvcc --version
nvidia-smi

The CUDA version shown by nvidia-smi describes driver support, not the loaded toolkit version. See NVIDIA’s CUDA compatibility guidance when using a newer toolkit with an older driver. Native GPU code and PTX have different compatibility requirements.

Compiling a simple CUDA program

Save the following example as gpu_square.cu. It launches a GPU kernel and checks the returned values:

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

void check(cudaError_t status) {
    if (status != cudaSuccess) {
        std::fprintf(stderr, "%s\n", cudaGetErrorString(status));
        std::exit(EXIT_FAILURE);
    }
}

__global__ void square(int *values) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < 256) values[i] = i * i;
}

int main() {
    int *values;
    check(cudaMallocManaged(&values, 256 * sizeof(int)));
    square<<<1, 256>>>(values);
    check(cudaGetLastError());
    check(cudaDeviceSynchronize());
    for (int i = 0; i < 256; ++i) {
        if (values[i] != i * i) return EXIT_FAILURE;
    }
    check(cudaFree(values));
    std::puts("GPU square check PASS (256 elements)");
    return EXIT_SUCCESS;
}

Compile it for an A100 and run it within a GPU allocation:

nvcc -O2 -arch=sm_80 gpu_square.cu -o gpu_square
./gpu_square

A successful run prints GPU square check PASS (256 elements). For code targeting multiple GPU architectures, see GPU Code Generation Options below.

Compiling the sample programs

You do not need to be using a GPU-enabled node to compile the sample programs but you do need at least one GPU to run them.

The following example batch submission script downloads the matching NVIDIA sample release, builds matrixMul with CMake, and runs its result check. It uses a new directory for each job, preserving previous results.

#!/bin/bash
#SBATCH --partition=gpu
#SBATCH --gres=gpu:a100:1
#SBATCH --ntasks=1
#SBATCH --cpus-per-task=4
#SBATCH --mem-per-cpu=2000
#SBATCH --time=00:20:00

set -e
export SLURM_EXPORT_ENV=ALL
module load CUDA/13.2.0 GCC/15.2.0 CMake/4.2.1-GCCcore-15.2.0
mkdir "cuda-example-$SLURM_JOB_ID"
cd "cuda-example-$SLURM_JOB_ID"
git clone --depth 1 --branch v13.2 https://github.com/NVIDIA/cuda-samples.git
cmake -S cuda-samples/Samples/0_Introduction/matrixMul -B matrix-build
cmake --build matrix-build --parallel "$SLURM_CPUS_PER_TASK"
srun ./matrix-build/matrixMul

The sample reports whether its matrix result matches its reference calculation. It is a demonstration rather than a production benchmark.

GPU Code Generation Options

To achieve the best possible performance whilst being portable, GPU code should be generated for the architecture(s) it will be executed upon.

This is controlled by specifying -gencode arguments to NVCC which, unlike the -arch and -code arguments, allows for ‘fatbinary’ executables that are optimised for multiple device architectures.

Each -gencode argument requires two values, the virtual architecture and real architecture, for use in NVCC’s two-stage compilation. For example, -gencode=arch=compute_80,code=sm_80 specifies a virtual architecture of compute_80 and real architecture sm_80.

To support future hardware of higher compute capability, an additional -gencode argument can be used to enable Just in Time (JIT) compilation of embedded intermediate PTX code. This argument should use the highest virtual architecture specified in other gencode arguments for both the arch and code i.e. -gencode=arch=compute_80,code=compute_80.

The minimum specified virtual architecture must be less than or equal to the Compute Capability of the GPU used to execute the code.

A100 GPUs have compute capability 8.0; H100 GPUs have compute capability 9.0. To build the example with native code for both and additional PTX:

nvcc -O2 gpu_square.cu -o gpu_square_portable \
   -gencode=arch=compute_80,code=sm_80 \
   -gencode=arch=compute_90,code=sm_90 \
   -gencode=arch=compute_90,code=compute_90

PTX JIT compilation requires compatible driver libraries, even when native code works through CUDA minor-version compatibility. NVIDIA’s forward-compatibility packages can provide these on supported GPU and driver combinations. If you encounter an unsupported PTX toolchain error, contact Research IT for advice on the appropriate CUDA module and compatibility configuration.

Further details of these compiler flags can be found in the NVCC Documentation, along with details of the supported virtual architectures and real architectures.

Documentation

Nsight Systems

Nsight Systems is a system-wide performance analysis tool designed to visualise an application’s algorithms and identify the largest opportunities to optimise.

A common use-case for Nsight Systems is to generate application timelines via the command line, which can later be visualised on a local computer using the GUI component. The Nsight Systems executable, nsys, is provided by loading a version of CUDA using a module file.

Note

Use a version of Nsight Systems that supports the CUDA toolkit used to compile your application.

To generate an application timeline with Nsight Systems CLI (nsys):

nsys profile --sample=none --cpuctxsw=none -o timeline ./gpu_square

Nsight systems can trace multiple APIs, such as CUDA and OpenACC. The --trace argument specifies which APIs should be traced. See the nsys profiling command switch options for further information.

nsys profile --sample=none --cpuctxsw=none -o timeline_cuda --trace cuda,nvtx,osrt,openacc ./gpu_square

Once this file has been downloaded to your local machine, it can be opened in nsys-ui via File > Open > timeline.nsys-rep

Nsight Compute

Note

Nsight Compute kernel profiling is currently unavailable on Stanage because access to GPU performance counters is disabled. It is expected to be enabled in future, but no availability date is confirmed. Nsight Systems timeline profiling remains available.

Profiling using nvprof

nvprof is not supported on the GPUs in Stanage (it does not support NVIDIA architectures >= SM80); please use Nsight Systems instead.

CUDA Training

The Research Software Engineering team have developed an undergraduate teaching module on CUDA; lecture notes and lecture recordings for that module are accessible here for anyone with a University account.

Determining the NVIDIA Driver version

Run the command:

cat /proc/driver/nvidia/version

Installation notes

CUDA was installed using Easybuild, build details can be found in $EBROOTCUDA/easybuild with a given module loaded.