CUDA (acronimo di Compute Unified Device Architecture) è un'architettura hardware per l'elaborazione parallela creata da NVIDIA. Tramite l'ambiente di sviluppo per CUDA, i programmatori di software possono scrivere applicazioni capaci di eseguire calcolo parallelo sulle GPU delle schede video NVIDIA.
I linguaggi di programmazione disponibili nell'ambiente di sviluppo CUDA sono estensioni dei linguaggi più diffusi per scrivere programmi. Il principale è 'CUDA-C' (C con estensioni NVIDIA), altri sono estensioni di Python, Fortran, Java e MATLAB.
Programmi che sfruttano l'architettura CUDA possono essere scritti anche utilizzando le librerie software OpenCL e DirectCompute.

CUDA dà accesso agli sviluppatori ad un set di istruzioni native per il calcolo parallelo di elementi delle GPU CUDA. Usando CUDA, le ultime GPU Nvidia diventano in effetti architetture aperte come le CPU. Diversamente dalle CPU, le GPU hanno un'architettura parallela con diversi core, ognuno capace di eseguire centinaia di processi simultaneamente: se un'applicazione è adatta per questo tipo di architettura, la GPU può offrire grandi prestazioni e benefici. Questo approccio alla risoluzione dei problemi è noto come GPGPU.

CUDA ha parecchi vantaggi rispetto alle tradizionali tecniche di calcolo sulle GPU che usano le API grafiche.

• Il codice può essere letto da posizioni arbitrarie in memoria.
• Memoria condivisa: CUDA espone una regione di 16kB di grandezza che può essere condivisa velocemente fra i thread. Questa può essere usata come una cache gestita dall'utente, rendendo disponibili grandi larghezze di banda che è possibile usare per strutture texture.
• Letture e scritture veloci, verso e dalla GPU.
• Supporto completo per divisioni intere e operazioni bit-a-bit, tra cui l'accesso a texture intere.

• CUDA era inizialmente un sottoinsieme del linguaggio C, privo di ricorsione e puntatori a funzione, più alcune semplici estensioni. Un singolo processo deve essere eseguito attraverso multiple disgiunzioni di spazi di memoria, diversamente da altri ambienti di runtime C. Le versioni più recenti tuttavia usano la sintassi C++. Il codice scritto per le versioni precedenti può non compilare, oppure può comportarsi in modo differente.
• Il rendering delle texture non è supportato.
• Per la doppia precisione, supportata a partire dalla serie GTX 2XX, ci sono diverse deviazioni dallo standard IEEE 754: l'arrotondamento al pari è l'unica approssimazione supportata per: reciproci, divisioni e radici quadrate. In singola precisione i NAN segnalati e denormalizzati non sono supportati; queste sono specifiche per istruzioni di base, rispetto ad una singola parola di controllo^{[non chiaro]}; e la precisione delle cifre decimali di divisioni o radici n-esime è molto minore rispetto alla singola precisione.
• La larghezza di banda e la latenza tra CPU e GPU può essere un collo di bottiglia.
• I thread devono essere eseguiti in multipli di 32 per ottenere migliori prestazioni, con un numero totale di thread nell'ordine di migliaia. I rami dei codici non influiscono nelle prestazioni, a condizione che ciascuno dei 32 thread prenda lo stesso cammino di esecuzione. Il modello di esecuzione SIMD diviene una limitazione significativa per diversi compiti, per esempio l'attraversamento di uno spazio partizionato di strutture dati durante il raytracing.
• Diversamente da OpenCL, GPU dotate di CUDA sono disponibili solo da NVIDIA (GeForce 8 serie superiori, Quadro e Tesla)
• La gestione delle eccezioni non è supportata nel codice CUDA a causa del degrado delle prestazioni risultante con migliaia di thread in esecuzione.
• Nelle versioni più recenti non esistono emulatori software per l'esecuzione del codice CUDA.

La piattaforma BOINC (Berkeley Open Infrastructure for Network Computing) ha diversi progetti di calcolo distribuito che supportano e sfruttano a pieno le potenzialità e i benefici della tecnologia CUDA.

I progetti BOINC supportati dalla piattaforma Nvidia CUDA sono:^[1]

• Collatz Conjecture: si pone come obiettivo la soluzione della celebre congettura di Collatz (o 3n + 1)
• DistrRTgen: è un progetto volto alla creazione di rainbow table distribuite gratuitamente
• Einstein@home: ha lo scopo di analizzare i dati prodotti da osservatori astronomici per rilevare delle onde gravitazionali da diverse fonti, come ad esempio i buchi neri
• GPUgrid.net: è un progetto nel campo della ricerca biomedica
• Milkyway@home: l'obiettivo di questo progetto è la creazione di un accurato modello tridimensionale della Via Lattea, utilizzando i dati forniti dallo Sloan Digital Sky Survey
• Moo!: permette di partecipare alla competizione RC5-72, è un wrapper di distributed.net
• PrimeGrid: esegue la ricerca dei numeri primi più elevati al mondo
• SETI@home: è il famosissimo progetto di calcolo distribuito, utilizza i dati del radiotelescopio di Arecibo per la ricerca di segnali correlati a intelligenze extraterrestri

Questo esempio di codice in C++ carica una texture da un'immagine dentro un array della GPU:

* cudaArray* cu_array;
texture<float, 2> tex;

// Allocate array
cudaChannelFormatDesc description = cudaCreateChannelDesc<float>();
cudaMallocArray(&cu_array, &description, width, height);

// Copy image data to array
cudaMemcpy(cu_array, image, width*height*sizeof(float), cudaMemcpyHostToDevice);

// Bind the array to the texture
cudaBindTextureToArray(tex, cu_array);

// Run kernel
dim3 blockDim(16, 16, 1);
dim3 gridDim(width / blockDim.x, height / blockDim.y, 1);
kernel<<< gridDim, blockDim, 0 >>>(d_odata, height, width);
cudaUnbindTexture(tex);

__global__ void kernel(float* odata, int height, int width)
{
   unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
   unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;
   float c = tex2D(tex, x, y);
   odata[y*width+x] = c;
}

Sotto un esempio in Python che computa il prodotto di due array nella GPU. Il linguaggio non ufficiale Python può essere ottenuto da PyCUDA.

import pycuda.driver as drv
import numpy
import pycuda.autoinit

mod = drv.SourceModule("""
__global__ void multiply_them(float *dest, float *a, float *b)
{
  const int i = threadIdx.x;
  dest[i] = a[i] * b[i];
}
""")

multiply_them = mod.get_function("multiply_them")

a = numpy.random.randn(400).astype(numpy.float32)
b = numpy.random.randn(400).astype(numpy.float32)

dest = numpy.zeros_like(a)
multiply_them(
        drv.Out(dest), drv.In(a), drv.In(b),
        block=(400,1,1))

print(dest-a*b)

Altri Python bindings per semplificare i prodotti tra matrici possono essere trovati su pycublas.

import numpy
from pycublas import CUBLASMatrix
A = CUBLASMatrix( numpy.mat([[1,2,3],[4,5,6]],numpy.float32) )
B = CUBLASMatrix( numpy.mat([[2,3],[4,5],[6,7]],numpy.float32) )
C = A*B
print C.np_mat()

• Calcolo distribuito
• GPU
• GPGPU
• Halide
• NVIDIA
• OpenCL

• Wikimedia Commons contiene immagini o altri file su CUDA

• Sito ufficiale, su developer.nvidia.com.
• (EN) Opere riguardanti CUDA (Computer architecture), su Open Library, Internet Archive.
• CUDA Community su Google+