Mandelbrot Cuda su Nvidia Jetson Nano

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;

}