Skip to content

Nvidia CUDA Compiler

Nvidia CUDA Compiler driver NVCCC is a proprietary, LLVM based, compiler by Nvidia intended for use with CUDA codes on both the CPU and GPU. NVCC separates these two parts and sends host code (the part of code which will be run on the CPU) to a C compiler (GCC, Intel C++ Compiler or Microsoft Visual C++ Compiler), and the device code (the part which will run on the GPU) to the GPU. The device code is further compiled by NVCC.


Following versions of CUDA are currently available:

You can load selected version as a module by following commands:

module load CUDA/11.3.1

You can load selected version as a module by following commands:

module load CUDA/11.4.1

You can load selected version as a module by following commands:

module load CUDA/12.0.0

Compiling with CUDA

CUDA code can be compiled directly on the login nodes and users do not have to compute nodes with GPU acceleration for the compilation. CUDA code can be compiled using the NVCC compiler. Thorough description of NVCC command options can be found on the Nvidia websites dedicated to CUDA compiler.

login01:~$ nvcc --version

NVCC comes with a large set of examples that cover both, elementary and intermediate, compilation codes. To test the compilation users should copy these examples to their home directory. To compile the code change directory to the particular example and run make command to start the compilation.

login01:~$ mkdir nvcc-tests
login01:~$ cp -r /storage-apps/easybuild/software/CUDA/<version>/samples/ ~/nvcc-tests/

login01:~$ cd nvcc-tests/1_Utilities/deviceQuery
login01:~/nvcc-tests/1_Utilities/deviceQuery$ make

To run the code, the user can use interactive slurm session to get access to a node with the Nvidia GPU and execute the binary file:

login01:~/nvcc-tests/1_Utilities/deviceQuery$ srun --job-name=nvcc-tests --partition=ngpu -G 1 --pty bash
login01:~/nvcc-tests/1_Utilities/deviceQuery$ ml CUDA

The expected output of the deviceQuery example executed on a node with a Nvidia SXM A100 40 GB is:

Cuda output
CUDA Device Query (Runtime API) version (CUDART static linking)

Detected 1 CUDA Capable device(s)

Device 0: "NVIDIA A100-SXM4-40GB"
  CUDA Driver Version / Runtime Version          12.0 / 11.4
  CUDA Capability Major/Minor version number:    8.0
  Total amount of global memory:                 40370 MBytes (42331013120 bytes)
  (108) Multiprocessors, (064) CUDA Cores/MP:    6912 CUDA Cores
  GPU Max Clock rate:                            1410 MHz (1.41 GHz)
  Memory Clock rate:                             1215 Mhz
  Memory Bus Width:                              5120-bit
  L2 Cache Size:                                 41943040 bytes
  Maximum Texture Dimension Size (x,y,z)         1D=(131072), 2D=(131072, 65536), 3D=(16384, 16384, 16384)
  Maximum Layered 1D Texture Size, (num) layers  1D=(32768), 2048 layers
  Maximum Layered 2D Texture Size, (num) layers  2D=(32768, 32768), 2048 layers
  Total amount of constant memory:               65536 bytes
  Total amount of shared memory per block:       49152 bytes
  Total shared memory per multiprocessor:        167936 bytes
  Total number of registers available per block: 65536
  Warp size:                                     32
  Maximum number of threads per multiprocessor:  2048
  Maximum number of threads per block:           1024
  Max dimension size of a thread block (x,y,z): (1024, 1024, 64)
  Max dimension size of a grid size    (x,y,z): (2147483647, 65535, 65535)
  Maximum memory pitch:                          2147483647 bytes
  Texture alignment:                             512 bytes
  Concurrent copy and kernel execution:          Yes with 5 copy engine(s)
  Run time limit on kernels:                     No
  Integrated GPU sharing Host Memory:            No
  Support host page-locked memory mapping:       Yes
  Alignment requirement for Surfaces:            Yes
  Device has ECC support:                        Enabled
  Device supports Unified Addressing (UVA):      Yes
  Device supports Managed Memory:                Yes
  Device supports Compute Preemption:            Yes
  Supports Cooperative Kernel Launch:            Yes
  Supports MultiDevice Co-op Kernel Launch:      Yes
  Device PCI Domain ID / Bus ID / location ID:   0 / 23 / 0
  Compute Mode:
   < Default (multiple host threads can use ::cudaSetDevice() with device simultaneously) >

deviceQuery, CUDA Driver = CUDART, CUDA Driver Version = 12.0, CUDA Runtime Version = 11.4, NumDevs = 1
Result = PASS

Code Example

Following is CUDA based code for two vector addition. To test it copy the code to and compile it with nvcc -o cuda-vector-addition. To execute the code open an interactive Slurm session on GPU accelerated node, as described above.

CUDA vector addition code
#define N (2048*2048)

#include <stdio.h>
#include <stdlib.h>

// Add vectors on GPU
__global__ void add_gpu(int *x, int *y, int *z, int n) {
    int index = threadIdx.x + blockIdx.x * blockDim.x;
    if (index < n)
        z[index] = x[index] + y[index];

// Add vectors on CPU
void add_cpu(int *x, int *y, int *z, int n) {
    for (int j = 0; j < n; j++)
        z[j] = x[j] + y[j];

// Generate vectors 
void random_ints(int *x, int n) {
    for (int j = 0; j < n; j++)
        x[j] = rand() % 10000; // random number between 0 and 9999

// Compare vectors on CPU 
int compare_ints(int *x, int *y, int n) {
    int pass = 0;
    for (int j = 0; j < N; j++) {
        if (x[j] != y[j]) {
            printf("Value differ at location %d, with values of %d and %d\n", j, x[j], y[j]);
            pass = 1;
    if (pass == 0)
    return pass;

int main(void) {

    int *x, *y, *z; // CPU copies x, y, z
    int *dev_x, *dev_y, *dev_z; // GPU copies x, y, z
    int size = N * sizeof(int); // Create N dimensional space for N integers

    // Allocate GPU copies of dev_x, dev_y, dev_z
    cudaMalloc((void**)&dev_x, size);
    cudaMalloc((void**)&dev_y, size);
    cudaMalloc((void**)&dev_z, size);

    // Allocate CPU copies of x, y, z
    x = (int*)malloc(size);
    y = (int*)malloc(size);
    z = (int*)malloc(size);

    // Create input vector with random integer numbers
    random_ints(x, N);
    random_ints(y, N);

    // Copy inputs to GPU
    cudaMemcpy(dev_x, x, size, cudaMemcpyHostToDevice);
    cudaMemcpy(dev_y, y, size, cudaMemcpyHostToDevice);

    // Launch add_gpu() kernel with blocks and threads
    add_gpu<<<N / THREADS_PER_BLOCK, THREADS_PER_BLOCK>>>(dev_x, dev_y, dev_z, N);

    // Copy GPU result back to CPU copy of z
    cudaMemcpy(z, dev_z, size, cudaMemcpyDeviceToHost);

    // Compare with CPU results
    int *z_h;
    z_h = (int*)malloc(size);
    add_cpu(x, y, z_h, N);
    compare_ints(z, z_h, N);

    // Clean CPU memory

    // Clean GPU memory

    return 0;