Ved rendering af 3D-grafik med et API som f.eks. OpenGL skriver man programmer direkte til GPU'en - de såkaldte shaders, som bl.a. har fået deres navn fra deres "oprindelige funktion": shading, altså til beregning af lys og skygger (se f.eks. Blinn-Phong shading).
Men eftersom GPU'er har en fundamentalt anderledes arkitektur end CPU'er, med mange tusinder af kerner, er der mulighed for at parallelisere sine beregninger på en helt anden skala end med CPU'er.
Dette førte til muligheden for at anvende f.eks. compute-shaders i OpenGL, som ikke er en del af en rendering pipeline, men simpelthen bare giver mulighed for at udnytte GPU-regnekraft.
Til applikationer, hvor der ikke er brug for at rendere grafik, er det dog muligt at skrive programmer til GPU'er uden at anvende shaders via f.eks. OpenGL. Det er her CUDA kommer ind i billedet!
Et simpelt program - vektor-addition
For at illustrere brugen af CUDA, vil vi se på et simpelt eksempel: vektor-addition.
Udgangspunktet er følgende C++-program:
#include <iostream>
#include <chrono>
#include <memory>
extern "C" void cuda_add_vectors(double* v1, double* v2, double* result, int len);
void add_vectors(double* v1, double* v2, double* result, int len)
{
for (auto i = 0; i < len; i++)
result[i] = v1[i] + v2[i];
}
template <typename T> double get_elapsed(T start, T end, int reps)
{
std::chrono::duration<double> elapsed = end - start;
auto ns = std::chrono::duration_cast<std::chrono::nanoseconds>(elapsed).count();
return static_cast<double>(ns) / (1000000.0 * reps);
}
int main()
{
const int repetitions = 100;
const int len = 50000000;
// Allocate vectors
std::shared_ptr<double> v1m(new double[len], std::default_delete<double[]>());
std::shared_ptr<double> v2m(new double[len], std::default_delete<double[]>());
std::shared_ptr<double> v3m(new double[len], std::default_delete<double[]>());
std::shared_ptr<double> v4m(new double[len], std::default_delete<double[]>());
double* v1 = v1m.get();
double* v2 = v2m.get();
double* result_cpu = v3m.get();
double* result_gpu = v4m.get();
// Fill vectors with some data
for (int i = 0; i < len; i++)
{
v1[i] = i*i;
v2[i] = 500.0 / (i + 1);
}
// ----------------------------------------------------
std::cout << "Running on CPU..." << std::endl;
auto start = std::chrono::system_clock::now();
for (auto i = 0; i < repetitions; i++)
add_vectors(v1, v2, result_cpu, len);
auto end = std::chrono::system_clock::now();
std::cout << "Done running on CPU. Elapsed time: " << get_elapsed(start, end, repetitions) << " ms." << std::endl;
// ----------------------------------------------------
// Note: First invocation has extra overhead
cuda_add_vectors(v1, v2, result_gpu, len);
std::cout << "Running on GPU..." << std::endl;
start = std::chrono::system_clock::now();
for (auto i = 0; i < repetitions; i++)
cuda_add_vectors(v1, v2, result_gpu, len);
end = std::chrono::system_clock::now();
std::cout << "Done running on GPU. Elapsed time: " << get_elapsed(start, end, repetitions) << " ms." << std::endl;
// ----------------------------------------------------
// Check that results are consistent
int bad_results = 0;
for (int i = 0; i < len; i++)
bad_results += (std::abs(result_cpu[i] - result_gpu[i]) > 1e-3) ? 1 : 0;
std::cout << "There are " << bad_results << " inconsistent results." << std::endl;
exit(0);
}
Der forberedes et par store vektorer, og deres sum beregnes først via funktionen add_vectors,
som bare laver beregningen på CPU'en via et almindeligt sekventielt for-loop
(man kan evt. speede det op med OpenMP eller lign. for bedre benchmarks af CPU vs. GPU performance).
Dernæst laves samme beregning på GPU via CUDA - koden til dette kommer vi til omlidt.
Bemærk at CUDA-koden ikke står direkte sammen med C++-koden, men skal kompileres separat,
og cuda_add_vectors-prototypen er angivet som extern "C" (for at undgå name mangling).
Bemærk også, at cuda_add_vectors køres en gang før selve benchmarken.
Dette er pga. ekstra overhead ved første kørsel af et program på GPU - så længe man bare kører samme GPU-program flere gange,
er der væsentligt mindre overhead.
Vektor-addition med CUDA
For at anvende CUDA direkte, skal man bruge NVIDIAs C-compiler, nvcc, som har en række udvidelser til C.
Udover almindelige C-funktioner stiller nvcc bl.a. funktionerne cudaMalloc,
cudaMemcpy og cudaFree til rådighed.
Disse funktioner fungerer stort set som de tilsvarende C-funktioner,
men bruges til at allokere og deallokere hukommelse på grafikkortet,
samt til at kopiere mellem main memory og grafikkortets hukommelse.
I CUDA skelnes der mellem host, som er CPU og main memory (alt andet end GPU), samt device, som er GPU'en.
Dette bruges bl.a. i cudaMemcpy hvor sidste parameter bl.a. kan være cudaMemcpyHostToDevice eller cudaMemcpyDeviceToHost,
som anvendes til at kopiere data hhv. fra main memory til GPU og den anden vej.
Vektor-addition via CUDA kan se således ud:
#include <stdio.h>
__global__ void cuda_run_add_vectors(double* v1, double* v2, double* result)
{
int index = threadIdx.x + blockDim.x * blockIdx.x;
result[index] = v1[index] + v2[index];
}
extern "C" void cuda_add_vectors(double* v1, double* v2, double* result, int len)
{
double* cv1;
double* cv2;
double* cresult;
int size = len * sizeof(double);
int blocks = 256;
int threadsPerBlock = ceil(len / (double)blocks);
int totalThreads = blocks * threadsPerBlock;
cudaMalloc(&cv1, totalThreads * sizeof(double));
cudaMalloc(&cv2, totalThreads * sizeof(double));
cudaMalloc(&cresult, totalThreads * sizeof(double));
cudaMemcpy(cv1, v1, size, cudaMemcpyHostToDevice);
cudaMemcpy(cv2, v2, size, cudaMemcpyHostToDevice);
cuda_run_add_vectors<<<threadsPerBlock, blocks>>>(cv1, cv2, cresult);
cudaMemcpy(result, cresult, size, cudaMemcpyDeviceToHost);
cudaFree(cv1);
cudaFree(cv2);
cudaFree(cresult);
}
Her er der to C-udvidelser i spil: Selve funktionen, som skal køre på GPU'en, er angivet som __global__.
Dette fortæller compileren, at denne funktion skal køres på GPU i stedet for CPU.
Derudover er der selve funktionskaldet af cuda_run_add_vectors.
Her anvendes <<<...>>> til at angive, at der skal startes en funktion på GPU'en,
og de to parametre angivet mellem <<< og >>> angiver hvor mange threads på GPU'en, som skal anvendes.
Threads på GPU'en er indelt er blokke, her 256 blokke, med lige mange threads i hver.
Så hvis der f.eks. er 128 threads i hver blok, vil der være 32768 threads i alt på GPU'en,
som kører cuda_run_add_vectors.
Hver instans af cuda_run_add_vectors, altså i hver thread, har adgang til nogle parametre,
som fortæller hvilken thread instansen kører; threadIdx.x, blockDim.x og blockIdx.x.
Disse anvendes til at beregne et indeks i vektoren, og bemærk at vektor-addition således udføres uden et loop,
men simpelthen ved at hvert indeks håndteres af et separat thread.
Kompilering og linking
Selve CUDA C-koden skal kompileres med nvcc mens C++-koden kan kompileres med f.eks. g++.
De resulterende object files skal linkes sammen, og der skal linkes til CUDA runtime.
Her er en simpel makefile som håndterer dette:
all: cudaexample cudaexample: cuda_add.o example.o g++ -L/usr/local/cuda/lib64 cuda_add.o example.o -o cudaexample -lcuda -lcudart example.o: example.cpp g++ -c example.cpp cuda_add.o: cuda_add.cu nvcc -c cuda_add.cu clean: rm -f *.o cudaexample
Bemærk at stien /usr/local/cuda/lib64 ikke nødvendigvis passer med dit setup.
Fejlfinding og profilering
Det er muligt at alt kompileres og linkes fint, men at der intet køres på GPU'en når du starter programmet. Dette kan f.eks. skyldes at CUDA ikke er installeret korrekt, eller hvis der ikke er en CUDA-kompatibel GPU til rådighed.
For at få information om fejl, kan følgende funktion bruges med CUDA C-koden:
void error_check()
{
cudaError_t code = cudaGetLastError();
if (code != cudaSuccess)
{
const char* error = cudaGetErrorString(code);
printf("[ERROR] Code: %d, message: %s\n", code, error);
}
}
Sørg for at kalde error_check lige efter kaldet til cuda_run_add_vectors for at få information om fejl.
Når det hele kører som det skal, kan du desuden bruge nvprof eller Nsight Compute/System til at profilere CUDA-programmer:
nvprof ./cudaexample
eller
ncu ./cudaexample ... eller ... nsys profile ./cudaexample
Dette giver en del information omkring f.eks. mængden af overhead ved brug af cudaMemcpy eller cudaMalloc.