NVIDIA CUDA Kompilátor¶
NVIDIA CUDA Compiler Driver (nvcc) je hlavný kompilátor
používaný na tvorbu CUDA aplikácií určených pre NVIDIA GPU. Je súčasťou
CUDA Toolkit a zodpovedá za kompiláciu GPU jadier (kernels) a za
koordináciu kompilácie hostiteľského kódu.
Kompilátor je založený na LLVM infraštruktúre a vo svojom vnútri rozdeľuje program na dva komponenty:
- Hostiteľský kód – kompilovaný štandardným CPU kompilátorom, napr. GCC alebo Intel C++ Compiler
- Device kód – kompilovaný pre architektúru GPU pomocou CUDA kompilačného pipeline
Tento dizajn umožňuje, aby sa CUDA programy integrovali so štandardnými C/C++ aplikáciami, pričom sa paralelné záťaže presúvajú na GPU.
Dostupné verzie CUDA¶
CUDA toolkity sú dostupné cez modulový systém, a dajú sa načítať priamo:
module load CUDA/<version>
Príklad:
module load CUDA/12.0.0
Zoznam všetkých dostupných CUDA modulov:
module avail CUDA
Overenie inštalácie CUDA¶
Po načítaní CUDA modulu overte, že CUDA kompilátor je dostupný:
nvcc --version
Typický výstup by mal obsahovať informácie o verzii CUDA toolkitu a o hostiteľskom kompilátore používanom NVCC.
Viditeľnosť GPU na výpočtových uzloch môžete overiť aj pomocou:
nvidia-smi
Kompilácia s CUDA¶
CUDA kód sa dá kompilovať priamo na prihlasovacích uzloch; na jeho kompiláciu nepotrebujete GPU výpočtový uzol.
CUDA kód sa dá kompilovať pomocou kompilátora NVCC. Podrobný popis možností príkazu NVCC nájdete na webe NVIDIA venovanom CUDA kompilátoru.
nvcc --version
NVCC prináša veľkú sadu príkladov pokrývajúcich základné
aj stredne pokročilé kompilačné úlohy. Na otestovanie kompilácie by
používatelia mali tieto príklady skopírovať do svojho domovského
adresára. Na kompiláciu prejdite do adresára konkrétneho príkladu
a spustite príkaz make, ktorý spustí kompiláciu.
mkdir nvcc-tests
cp -r /storage-apps/easybuild/software/CUDA/samples/ nvcc-tests/.
cd nvcc-tests/Samples/1_Utilities/deviceQuery
make
Spúšťanie CUDA programov¶
CUDA programy musia byť spustené na GPU uzloch. Interaktívna Slurm relácia sa môže spustiť pomocou:
srun --job-name=nvcc-tests --partition=gpu -G 1 --pty bash # Ran on login node
ml CUDA # Ran on gpu accelerated node
./deviceQuery # Ran on gpu accelerated node
exit # Ran on gpu accelerated node
Očakávaný výstup príkladu deviceQuery spusteného na uzle s GPU NVIDIA A100-SXM4-40GB je:
Devana 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
Príklad kódu¶
Nasleduje CUDA kód na sčítanie vektorov. Na otestovanie skopírujte kód
do cuda-vector-addition.cu a skompilujte ho pomocou nvcc.
CUDA vector addition code
#define N (2048*2048)
#define THREADS_PER_BLOCK 512
#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)
printf("Success\n");
else
printf("Fail\n");
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
free(x);
free(y);
free(z);
free(z_h);
// Clean GPU memory
cudaFree(dev_x);
cudaFree(dev_y);
cudaFree(dev_z);
return 0;
}
Na spustenie kódu ho skompilujte na prihlasovacom uzle a potom otvorte interaktívnu Slurm reláciu na uzle s GPU akceleráciou:
ml CUDA/12.0.0
nvcc cuda-vector-addition.cu -o cuda-vector-addition
srun --job-name=nvcc-tests --partition=gpu -G 1 --pty bash # Ran on login node
Spustenie na uzle s GPU akceleráciou:
ml CUDA/12.0.0 # Ran on gpu accelerated node
./cuda-vector-addition # Ran on gpu accelerated node
Očakávaná výstupná správa je Success.
Odporúčané postupy vývoja CUDA¶
Pri vývoji CUDA aplikácií na HPC systémoch:
- Vždy kompilujte pomocou CUDA modulu
- Spúšťajte GPU programy iba v GPU frontách
- Zlaďte verziu CUDA s kompatibilitou ovládača (driveru)
- Používajte MPI + CUDA pre multi-GPU aplikácie
- Testujte s malými vstupmi pred rozsiahlymi spusteniami