CUDA (acrònim de Compute Unified Device Architecture (Arquitectura de còmput de dispositius unificats)) és una plataforma de computació paral·lela i model d'Interfície de programació d'aplicacions (API) creada per Nvidia per permetre a desenvolupadors i enginyers de programari accelerar l'execució dels seus codis fent servir Unitats de processament gràfic (GPU) amb capacitat CUDA per a processament de caràcter general.
La plataforma CUDA és una capa de software que dona accés directe al conjunt virtual d'instruccions de la GPU i als seus elements de comput paral·lel a efectes d'executar nuclis de cómput (CUDA Kernels).
La plataforma va ser desenvolupada amb l'objectiu de treballar conjuntament amb llenguatges de programació com C, C++ i Fortran, tot i que posteriorment ha acabat implementant-se en un ampli espectre de llenguatges de programació com Java, Lua, MATLAB… Aquesta accessibilitat facilita als especialistes de programació paral·lela l'ús dels recursos que ofereix una GPU, en contrast amb altres APIs natives pròpiament de la prográmacio de gráfics (com Direct3D o OpenGL), les quals tenien un enfóc purament gráfic i requerien grans coneixements en aquest ambit.
CUDA dona suport a entorns de treball (Frameworks) orientats a la programació amb acceleradors com OpenACC i OpenCL.
Història 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ó.
Aquest 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.
Jerarquia de fils d'execució 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. alt=Jerarquia de fils d'execució en CUDA|miniatura|Jerarquia simplificada de fils d'execució en CUDA. 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ó:miniatura|Organització detallada del fils d'execució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).
Jerarquia de memòria miniatura|CUDA memòria|253x253pxEls 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ó.
Accés a la memòria del dispositiu Memòria global 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.
Memòria compartida La memòria compartida està ubicada en el xip obtenint una 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.
Memòria local 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).
Memòria de textures i superfícies 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.
Estructura d'una aplicació CUDA 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.
Creació de kernels en CUDA CUDA Kernel CUDA 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ó.
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; }
Invocacions a un kernel 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);
Exemple 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]; }