CUDA

La unitat de processament gràfic (GPU), és un processador especialitzat, dirigit en els seus inicis a les demandes de càlcul intensiu de gràfics d'alta resolució en 3D en temps real. A l'any 2012, les GPU havien evolucionat fins haberse convertit en sistemes multi-core masius, altament paral·lels i enfocats a la manipulació eficient de grans blocs de dades. Aquest disseny és més eficaç que la unitat central de processament de propòsit general (CPU) per als algorismes que han de treballar grans blocs de dades i requerint poca interacció.

Aquesta nou paradigma va posar de manifest la gran necessitat d'un llenguatge de programació no orientat estrictament als gràfics que permetés la utilitzar totes les prestacions d'aquests dispositius d'una forma versatil i de relatiu alt nivell, assentant la importancia del llenguatge de programació CUDA present al mercat des de 2007.

Per l'execució de programes paral·lels a través de kernels CUDA Kernels s'utilitza una jerarquia de fils que conté tres categories anidades: Thread, ThreadBlock i Grid. Aquests estan anidats com mostra imatge dreta, on podem veure: En blau la representació dels threads, englobats en els blocks grisos, dintre d'un grid, representat en verd.

Com aquí s'observa, la quantitat de threads i blocks creats ve definida des de la crida a la funció. La següent imatge, mostra de forma més detallada com després aquests fils s'organitzen, així com una serie d'identificadors associats. Aquests identificadors tenen una traducció directa en variables accesibles del llenguatge, que descriuen el context de treball en termes de dimensions i posició:

CUDA introdueix dos tipus de dades per aquestes descripcions:

• unit3: Vector interconstruït de tipus sencer. Funció constructor: unit3 nom_variable(x,i,z).
• dim3: Vector interconstruït i basat en unit3. Utilitzat per especificar dimensions. Funció constructor: dim3 nom_variable(x,i,z) (qualsevol component no especificat s'inicialitza a 1).

Amb aquest tipus de dades construeix 4 variables, 2 referents a l'identificador de localització i 2 referents a la mida del context:

• ThreadIdx: Variable de tipus unit3 que conté l'índex d'un thread (dins d'un Block).
• BlockIdx: Variables de tipus unit3 que conté l'índex d'un block (dins d'un Grid).
• BlockDim: Variable de tipus dim3 que conté la dimensió d'un Block.
• GridDim: Variable de tipus dim3 que conté la dimensió d'una Grid.

Els fils poden estar organitzats, d'acord amb l'aplicació, en blocs d'una dimensió (per manipular vectors), de dues dimensions (per manipular matrius) i de tres dimensions (per manipular arranjaments tridimensionals).

Els threads de CUDA poden accedir a múltiples memòries durant l'execució que es poden dividir en aquests tres casos:

• Memòria privada per thread d'ús exclusiu
• Memòria per bloc, és visible per tots els threads del bloc i té el mateix temps de vida que el bloc
• Memòria global accessible per tots els threads

Hi ha també dos memòries addicionals optimitzades per diferents usos accessibles per tots els threads:

• Memòria constant
• Memòria de textures i superfícies

Els espais de memòria global, constant i de textura són persistents entre kernels llençats en la mateixa aplicació.

La memòria global es troba a la memòria del dispositiu, l'accés és fa a través de transaccions de memòria de 32, 64 o 128 bytes. Aquestes operacions de memòria han d'estar alineades de manera natural: només es poden llegir o escriure a memòria segments de 32, 64 o 128 bytes del dispositiu alineats a la seva mida (és a dir, la primera adreça és un múltiple de la seva mida).

Quan un bloc executa una instrucció que accedeix a la memòria global, converteix els accessos de memòria dels fils dins d'un bloc en una o més d'aquestes transaccions de memòria depenent de la mida de la paraula a la qual s'accedeixi en cada fil i la distribució de les adreces de memòria a través dels fils. En general, com més transaccions siguin necessàries, es transferiran més paraules no utilitzades, a més de les paraules que accedeixen als fils, reduint-ne el rendiment. Per exemple, si es genera una transacció de memòria de 32 bytes per a l'accés de 4 bytes de cada fil, el rendiment es divideix en 8.

La memòria compartida està ubicada en el xip obtenint un amplada de banda més gran i una latència menor que la memòria local o global.

Per aconseguir una amplada de banda major la memòria compartida es divideix en mòduls de memòria de la mateixa mida anomenats bancs als quals es poden accedir de forma simultània. Qualsevol sol·licitud de lectura o escriptura a memòria feta en N direccions que cauen en N bancs de memòria diferents, és a dir, podran ser atesos simultàniament donant com a resultat una amplada de banda global que és n vegades més alt que l'amplada de banda d'un sol mòdul.

Tanmateix, si dues adreces d'una sol·licitud de memòria cauen al mateix banc de memòria, hi ha un conflicte de banc i l'accés ha de ser serialitzat. El maquinari divideix una sol·licitud de memòria amb conflictes de banc en tantes sol·licituds com sigui necessari, disminuint el rendiment per un factor igual al nombre de sol·licituds de memòria separades. Si el nombre de sol·licituds de memòria separades és n, es diu que la sol·licitud de memòria inicial provoca conflictes de banc en forma de n.

L'espai de memòria local resideix en la memòria del dispositiu, de manera que els accessos a la memòria local tenen la mateixa latència i amplada de banda baixos que els accessos de memòria global i estan subjectes als mateixos requisits de coalescència de memòria que es descriu als accessos de memòria del dispositiu. No obstant això, la memòria local s'organitza de manera que s'accedeix a paraules de 32 bits consecutives mitjançant identificacions de fil consecutives. Per tant, els accessos s'integraran íntegrament sempre que tots els fils d'un warp accedeixin a la mateixa adreça relativa (per exemple, el mateix índex en una variable de matriu, el mateix membre en una variable d'estructura).

Els espais de memòria per textures i superfícies resideixen en la memòria del dispositiu. El cache de textura està optimitzada per a localitats espacials 2D, cada thread del mateix bloc que llegeix una textura o superfície aconsegueixin un millor rendiment.

Les aplicacions en CUDA estan pensades per executar codi en forma vectorial aprofitant l'explotació del paral·lelisme al màxim. Una aplicació CUDA pot estar composta per parts de codi C seqüencials executades en CPU i kernels paral·lels executats en GPU combinant les dues infraestructures.

#include <cuda.h>void vecAdd(float* A, float* B, float* C, int n){ int size = n* sizeof(float); float* A_d, B_d, C_d; … 1. Reservem en la memoria del dispositus els arrays A, B i C cudaMalloc((void **) &A_d, size); cudaMalloc((void **) &B_d, size); cudaMalloc((void **) &C_d, size); // Copiem A i B a la memoria del dispositiu cudaMemcpy(A_d, A, size, cudaMemcpyHostToDevice); cudaMemcpy(B_d, B, size, cudaMemcpyHostToDevice); 2. // Executem el kernel de codi en el dispositiu per tal d'executar la suma de vectors 3. // Copiem el resultat de C del dispositiu al host (CPU) cudaMemcpy(C, C_d, size, cudaMemcpyDeviceToHost); // Lliurem la memoria utilitzada al dispositiu cudaFree(A_d); cudaFree(B_d); cudaFree (C_d);}

Per tal de fer les operacions dites anteriorment comptem amb les següents funcions implementades en la llibreria "cuda.h":

• cudaMalloc(Direcció del punter, Mida en bytes): Per reservar memòria al dispositiu.
• cudaFree(Punter a les dades): Per alliberar memòria prèviament reservada.
• cudaMemcpy(Punter de destí, Punter d'origen, Mida en bytes, Direcció de la transferència): Copia de blocs d'informació d'un punter a un altre.
• cudaMemset(Punter, valor, Mida en bytes): Inicialitza una quantitat de bytes a un valor concret.

La direcció de les dades per la funció cudaMemcpy pot ser una de les següents constants:

• cudaMemcpyHostToHost: Copia dades de memòria principal a memòria principal.
• cudaMemcpyHostToDevice : Copia dades de memòria principal a la memòria del dispositiu.
• cudaMemcpyDeviceToHost : Copia dades del dispositiu a memòria principal.
• cudaMemcpyDeviceToDevice : Copia dades del dispositiu al dispositiu.
• cudaMemcpyDefault: Es basa en els punters per decidir a on es copien les dades.

CUDA KernelCUDA incorpora un nou model d'execució diferent al model seqüencial de les computadores tradicionals. A CUDA, el codi serà executat per múltiples fils d'execció a la vegada (centenars o milers). La programació ha de ser modelada segons la jerarquia de fils d'execució.^[7]

Una funció kernel ha de ser invocada per la CPU (host), però s'executa exclusivament en GPU.

Els kernels tenen dues característiques fonamentals:

• Els kernels no poden retornar un valor de forma explícita. Els resultats han de ser desats sobre estructures de dades que s'hagin enviat a la funció en el moment de la crida.
• Els kernels generen de forma explícita la jerarquia de fils d'execció quan es realitza la crida segons el nombre de blocs de fils d'execució i el nombre de fils d'execució per bloc. Així doncs, una mateixa funció CUDA pot ser cridada diverses vegades en un mateix codi, i tenir una jerarquia de fils d'execució diferents segons la crida.

// Declaració del kernel CUDA.__global__ void foo(double *vector,int limit){ // Calcular l'identificador del thread int identificador = blockIdx.x*blockDim.x + threadIdx.x; // Comprovar que estem dintre del rang del vector if (identificador < limit) vector[identificador] = 0;}

Una vegada declarada la funció, és recomanable declarar variables per a gestionar la jerarquia de fils d'execució.

 int CTA_NUM = (N + 1023)/1024; dim3 blocks(CTA_NUM, 1, 1); dim3 threadsPerBlock (1024, 1, 1);

Finalment, es realitza la crida al kernel

 foo<<< blocks, threadsPerBlock >>>(d_vector_in, d_vector_out, N);

Codi d'exemple: Suma de vectors

#include <stdio.h>#include <stdlib.h>#include <math.h>// Declaració del kernel CUDA.__global__ void sumaVectors(double *a, double *b, double *c, int n){ // Calcular l'identificador del thread int id = blockIdx.x*blockDim.x+threadIdx.x; // Comprobar que estem dintre del rang del vector if (id < n) c[id] = a[id] + b[id];}// Funció principalint main(int argc, char* argv[]){ // Mida del vector int n = 1024*1024; // Declarem els punters d'entrada per la memòria del host  double *host_vector_a; double *host_vector_b; // Declarem els punters de sortida per la memòria del host  double *host_vector_c; // Declarem els punters d'entrada per la memòria del dispositiu (GPU) double *device_vector_a; double *device_vector_b; // Declarem els punters de sortida per la memòria del dispositiu (GPU) double *device_vector_c; // Calculem la mida del vector en bytes size_t bytes = n*sizeof(double); // Reservar memòria a l'espai del host per als vectors host_vector_a = (double*)malloc(bytes); host_vector_b = (double*)malloc(bytes); host_vector_c = (double*)malloc(bytes); // Reservar memòria a l'espai del dispositiu (GPU) per als vectors cudaMalloc(&device_vector_a, bytes); cudaMalloc(&device_vector_b, bytes); cudaMalloc(&device_vector_c, bytes); int i; // Inicialitzar valors del vectors for(i = 0; i < n; i++) { host_vector_a[i] = sin(i)*sin(i); host_vector_b[i] = cos(i)*cos(i); } // Copiar els vectors del host al dispositiu (D <-- H) cudaMemcpy(device_vector_a, host_vector_a, bytes, cudaMemcpyHostToDevice); cudaMemcpy(device_vector_b, host_vector_b, bytes, cudaMemcpyHostToDevice); int blockSize, gridSize; // Nombre de threads per bloc blockSize = 1024; // Nombre de blocks al grid gridSize = (int)ceil((float)n/blockSize); // Executem el kernel sumaVectors<<<gridSize, blockSize>>>(device_vector_a, device_vector_b, device_vector_c, n); // Copiem el vector de sortida del dispositiu al host (H <-- D) cudaMemcpy(host_vector_c, device_vector_c, bytes, cudaMemcpyDeviceToHost); // Calculem el resultat final double sum = 0; for(i=0; i<n; i++) sum += host_vector_c[i]; printf("final result: %f\n", sum/n); // Alliberem la memòria del dispositiu cudaFree(device_vector_a); cudaFree(device_vector_b); cudaFree(device_vector_c); // Alliberem la memòria del host free(host_vector_a); free(host_vector_b); free(host_vector_c); return 0;}

Codi exemple: Multiplicació per 2 de tots els pixels en una determinada zona

__global__ void PictureKernel(float* d_Pin, float* d_Pout, int n,int m){ // Calculem la fila de d_Pin i d_Pout que estem processant. int Row = blockIdx.y*blockDim.y + threadIdx.y; // Calculem la columna de d_Pin i d_Pout que estem processant. int Col = blockIdx.x*blockDim.x + threadIdx.x; // Cada thread computarà 1 element de d_Pin generant el corresponent d_Pout if ((Row < m) && (Col < n)) { d_Pout[Row*n+Col] = 2*d_Pin[Row*n+Col]; }}

Com els diferents fils col·laboren entre ells i poden compartir dades, es requereixen unes directives de sincronització. En un kernel, es pot declarar una barrera afegint una crida a __syncthreads(), amb això tots els fils s'esperaran fins que tots els fils arribin al mateix punt.

^[9]

Nota: Totes les línies que falten o les entrades buides reflecteixen alguna falta d'informació sobre aquest article exacte..
^[10]

^[11]

^[16]

Per obtenir més informació, vegeu l'article: «NVIDIA CUDA Compute Capability Comparative Table». and read Nvidia CUDA programming guide.^[17]

Nivell de CUDA soportat per la tarjeta GPU. Vegeu també a: Nvidia:

• CUDA SDK 1.0 suport per capacitat de cómput 1.0 – 1.1 (Tesla)^[18]
• CUDA SDK 1.1 suport per capacitat de cómput 1.0 – 1.1+x (Tesla)
• CUDA SDK 2.0 suport per capacitat de cómput 1.0 – 1.1+x (Tesla)
• CUDA SDK 2.1 – 2.3.1suport per capacitat de cómput 1.0 – 1.3 (Tesla)^{[19][20][21][22]}
• CUDA SDK 3.0 – 3.1 suport per capacitat de cómput 1.0 – 2.0 (Tesla, Fermi)^[23][24]
• CUDA SDK 3.2 suport per capacitat de cómput 1.0 – 2.1 (Tesla, Fermi)^[25]
• CUDA SDK 4.0 – 4.2 suport per capacitat de cómput 1.0 – 2.1+x (Tesla, Fermi, more?)
• CUDA SDK 5.0 – 5.5 suport per capacitat de cómput 1.0 – 2.1+x (Tesla, Fermi, more?)
• CUDA SDK 6.0 suport per capacitat de cómput 1.0 – 3.5 (Tesla, Fermi, Kepler)
• CUDA SDK 6.5 suport per capacitat de cómput 1.1 – 5.x (Tesla, Fermi, Kepler, Maxwell). Last version with support for compute capability 1.x (Tesla)
• CUDA SDK 7.0 – 7.5 suport per capacitat de cómput 2.0 – 5.x (Fermi, Kepler, Maxwell)
• CUDA SDK 8.0 suport per capacitat de cómput 2.0 – 6.x (Fermi, Kepler, Maxwell, Pascal). Last version with support for compute capability 2.x (Fermi)
• CUDA SDK 9.0 – 9.2 suport per capacitat de cómput 3.0 – 7.2 (Kepler, Maxwell, Pascal, Volta)
• CUDA SDK 10.0 – 10.1 suport per capacitat de cómput 3.0 – 7.5 (Kepler, Maxwell, Pascal, Volta, Turing)

^[26]

'*' – OEM - Només productes.

• Common Lisp – cl-cuda
• Clojure – ClojureCUDA
• Fortran – FORTRAN CUDA Arxivat 2012-04-10 a Wayback Machine., PGI CUDA Fortran Compiler
• F# – Alea.CUDA
• Haskell – Data.Array.Accelerate
• IDL – GPULib Arxivat 2012-11-07 a Wayback Machine.
• Java – jCUDA, JCuda, JCublas, JCufft, CUDA4J
• Julia – CUDAnative.jl Arxivat 2019-05-05 a Wayback Machine.^[30]
• Lua – KappaCUDA
• Mathematica – CUDALink
• MATLAB – Parallel Computing Toolbox, MATLAB Distributed Computing Server,^[31] and 3rd party packages like Jacket.
• .NET – CUDA.NET, Managed CUDA, CUDAfy.NET .NET kernel and host code, CURAND, CUBLAS, CUFFT
• Perl – KappaCUDA, CUDA::Minimal, AI::MXNet::CudaKernel
• Python – Numba, NumbaPro, PyCUDA, KappaCUDA, Theano Arxivat 2020-11-08 a Wayback Machine.
• Ruby – KappaCUDA (Broken link)
• R – gpuRcuda