Ho ripreso in mano la scheda Jetson Nano per iniziare a programmare in CUDA
Gli esempi di CUDA hanno gia' un sorgente per l'insieme di Mandelbrot ma visto che era a scopo didattico sono partito da zero
L'idea e' quella di usare la GPU per i calcoli. Ogni thread calcola una riga dell'immagine. Visto che ogni block della Nvidia puo' gestire al massimo 1024 thread la massima dimensione che il sorgente puo' generare e' 1024x1024
Per condividere una memoria tra la GPU e la CPU (ognuna puo' accere solo alle proprie risorse) si usa cudaMallocManaged. CudaMemset port i valori della matrice tutti a zero
la funzione kernel e'e quella che viene eseguita dalla GPU. Una volta lanciato il kernel la GPU restituisce il controllo alla CPU solo quando tutti i threads sono terminati. La variabile threadIdx viene utilizzate come indice di colonna dell'immagine
al termine l'array viene convertito in una immagine ppm a scala di grigio (nello specifico si sono solo due colori)
per compilare si usa nvcc con gli switch per indicano la compilazione sull'architettura 50 (relativa alla GPU sulla Jetson) con l'ottimizzazione fast_math
nvcc -use_fast_math -arch=sm_50 mand.cu -o mand
Ho usato VSCode con il terminale ma per programmare Cuda si puo' usare NSight (una versione di Eclipse con plugin NVidia) contenuto nei Cuda Tools
questo il codice
#include <stdio.h>
#include <time.h>
__global__ void kernel (int max, char *matrice)
{
int id = threadIdx.x;
/*zoom1
const double CxMin=-0.3041;
const double CxMax=-0.1874;
const double CyMin=-0.8867;
const double CyMax=-0.7699;*/
/*zoom2*/
const double CxMin=-0.2;
const double CxMax=-0.074;
const double CyMin=-1.058;
const double CyMax=-0.933;
/*insieme completo
const double CxMin=-2.5;
const double CxMax=1.5;
const double CyMin=-2.0;
const double CyMax=2.0;*/
const int iterazioni = 4096;
double x_new,y_new,x,y;
double a,b;
double PixelWidth=(CxMax-CxMin)/max;
double PixelHeight=(CyMax-CyMin)/max;
for (int s=0;s<max;s++)
{
x = 0;
y = 0;
a = CxMin + (PixelWidth*id);
b = CyMin + (PixelHeight*s);
for (int k=0;k<iterazioni;k++)
{
x_new = (x*x)-(y*y)+a;
y_new = (2*x*y)+b;
if (((x_new*x_new)+(y_new*y_new))>4)
{
// colora il punto
matrice[id*max+s] =(k%2)*255;
k=iterazioni;
}
x = x_new;
y = y_new;
}
}
}
int main(void)
{
clock_t tic = clock();
// crea la matrice in GPU
int dimensione = 1000;
char *matcuda;
// crea un array Cuda
cudaMallocManaged(&matcuda,dimensione*dimensione*sizeof(char));
// azzera tutti i valori dell'array
cudaMemset(matcuda,0,dimensione*dimensione*sizeof(char));
// 1024 e' il numero di threads che vengono eseguiti in contemporanea
// e' il valore massimo per ogni block della NVidia
// tra parentesi si possono passare variabili al kernel
kernel <<<1, dimensione>>>(dimensione,matcuda);
cudaDeviceSynchronize();
clock_t toc = clock();
/*
for (int s=0;s<dimensione;s++)
{
for (int t=0;t<dimensione;t++)
{
printf("%i",matcuda[(dimensione*s)+t]);
}
printf("\n");
}
*/
printf("Elapsed: %f seconds\n", (double)(toc - tic) / CLOCKS_PER_SEC);
FILE *fpi = fopen("mand.ppm", "wb");
fprintf(fpi, "P5\n%d %d\n255\n", dimensione, dimensione);
fwrite(matcuda, 1, dimensione*dimensione, fpi);
fclose(fpi);
cudaFree(matcuda);
return 0;
}
