Cuda Mandelbrot

 

 

 

#include <stdio.h>
//questa libreria crea png senza altre dipendenze
#include "libattopng.h"

#define W  5000 //width image
#define H  5000 //height image
#define TX 32 // number of threads per block along x-axis
#define TY 32 // number of threads per block along y-axis

__global__
void Kernel(int *d_out, int w, int h, int itera)
{
    //nel kernel viene lanciato un thread per ogni pixel dell'immagine
    //per calcolare la riga e la colonna si usa la formula sottostante
    int c = blockIdx.x*blockDim.x + threadIdx.x; // colonna
    int r = blockIdx.y*blockDim.y + threadIdx.y; // riga
    int i = r*w + c; // posizione del pixel nell'array lineare
    if ((c >= w) || (r >= h)) return;

    // inizia il ca
    const float x_min = -0.727;
    const float x_max = -0.730;
    const float y_min = -0.247;
    const float y_max = -0.250;

    float dx,dy;

    dx = (x_max-x_min)/w;
    dy = (y_max-y_min)/h;


    float k = x_min+(r*dx);
    float j = y_min+(c*dy);
    d_out[i] = 0;

    double x,y,x_new, y_new = 0.0;
    for (int f=0; f<itera;f++)
            {
                    x_new = (x*x) - (y*y) + k;
                    y_new = (2*x*y) +j;
                    if (((x_new*x_new)+(y_new*y_new))>4)
                    {
                        d_out[i] = f;
                        return;
                    }
                    x = x_new;
                    y = y_new;
            }
 }

int main()
{
 // alloca e mette a zero (malloc alloca e basta) un array di int di dimensione u
 // uguale all'immagine desiderata sulla ram della CPU
 int *out = (int*)calloc(W*H, sizeof(int));

 // alloca un array nella GPU delle stesse dimensioni
 // la GPU viene definita come device mentre la CPU e' l'host
 int *d_out; // pointer for device array
 cudaMalloc(&d_out, W*H*sizeof(int));

 // definisce i threads
 const dim3 blockSize(TX, TY);
 const int bx = (W + TX - 1)/TX;
 const int by = (W + TY - 1)/TY;
 const dim3 gridSize = dim3(bx, by);
 //lancia il kernel sulla GPU passando come parametro l'array sulla GPU
 Kernel<<<gridSize, blockSize>>>(d_out, W, H,4096);

 //trasferisce il contenuto dell'array dalla GPU alla RAM della CPU
 cudaMemcpy(out, d_out, W*H*sizeof(int), cudaMemcpyDeviceToHost);

 //crea una PNG partendo dall'array
 libattopng_t* png = libattopng_new(W, H, PNG_GRAYSCALE);
 
 int x, y;
 for (y = 0; y < W; y++) {
   for (x = 0; x < H; x++) {
     libattopng_set_pixel(png, y, x, out[x+y*H]%255);
   }
 }
 libattopng_save(png, "mandelbrot.png");
 libattopng_destroy(png);


 cudaFree(d_out);
 free(out);
 return 0;
 }
 

 

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;

}

Primi passi con NVidia Jetson Nano

L'impostazione di base e' stata ridotta al minimo. Il sistema e' basato sulla scheda Jetson Nano, una SD Card da 64 Gb, un dongle USB Wifi Edimax N150, una tastiera Logitech K400 Plus, un alimentatore esterno da 5V 4A

Per il case ho preso questo progetto  https://www.thingiverse.com/thing:3518410
E' piuttosto delicato su alcuni dettagli ma in generale va bene

Il primo problema riguarda l'alimentazione. Visto che non mi fidavo di una alimentazione via USB visti i  consumi della scheda ho utilizzato un alimentatore esterno. Il problema e' che per abilitare l'alimentazione esterna e' necessario chiudere con un Jumper J48...ed il jumper non e' incluso. Sono impazzito a trovare un jumper a giro per casa ed alla fine ho dovuto cannabalizzare un vecchio HD PATA (il jumper serviva a selezionare salve and master). La Jetson ha due modalita' di utilizzo 5W e 10 W (di default e' la seconda)

Da linea di comando si puo' modificare con (0=10W, 1=5W)

$ sudo nvpmodel -m 0 

$ sudo jetson_clocks

Nonostante sia sconsigliato di usare un convertitore da HDMI a DVI io ho lo ho utilizzato con successo per la connessione ad vecchio monitor

Invece mi ha fatto impazzire la Edimax che continuava a disconettersi dall'hotpost. La soluzione e' stata mettere in blacklist il modulo con

echo "blacklist rtl8192cu" | sudo tee -a /etc/modprobe.d/blacklist.conf

Si puo' avere una modalita' headless semplicemente collegando un cavo microUSB e puntando il browser del portatile su 192.168.55.1:8888 e password dlinano

In ogni caso il server SSH e' gia' installato e disponibile (ho collegato un cavo cross con un indirizzo fisso per la rete via cavo per estendere la modalita' headless)

Per provare CUDA Jetson ha gia' tutto installato. Si deve pero' settare la Path e copiare gli esempi nella propria home

export PATH=${PATH}:/usr/local/cuda/bin 
export LD_LIBRARY_PATH=${LD_LIBRARY_PATH}:/usr/local/cuda/lib64

Per gli esempi di CUDA

/usr/local/cuda-10.0/bin/cuda-install-samples-10.0.sh . (attenzione al punto finale)

ATTENZIONE : operazione lenta in compilazione

Jetson Nano 4 Giga di Ram e puo' essere poca per reti neurali impegnativi. Si puo' aumentare usando uno swap file

$ sudo fallocate -l 8G /mnt/4GB.swap 

$ sudo mkswap /mnt/4GB.swap 

$ sudo swapon /mnt/4GB.swap

e modificando il file /etc/fstab aggiungendo

/mnt/4GB.swap none swap sw 0 0

Ultima cosa: di default l'uscita audio e' impostata su Analogico e non su HDMI.

link utile

https://github.com/jkjung-avt/jetson_nano.git

Cuda Toolkit 8.0 su Ubuntu 18.04 LTS

Ho portatile T430 con scheda NVS 5400 M, una scheda con CUDA Capability 2.1 e Codename Fermi. Il supporto per queste schede e' terminato con Cuda Toolkit 8.0 e quindi non e' banale installare l'ambiente di sviluppo







Per prima cosa ovviamente si devono avere installati i driver proprietari di NVidia da Software&Updates

A questo punto si deve scaricare Cuda ToolKit 8, l'ultimo compatibile con Fermi ma si deve effettuare anche un downgrade di gcc (Cuda Toolkit 8 era compatibile con Ubuntu 16.04(

sudo apt-get install gcc-5 g++-5
sudo update-alternatives --install /usr/bin/gcc gcc /usr/bin/gcc-5 70
sudo update-alternatives --install /usr/bin/g++ g++ /usr/bin/g++-5 70






sh ./cuda_8.0.61_375.26_linux.run --tar mxvf
sudo apt install libmodule-install-perl
sudo cp InstallUtils.pm /usr/lib/x86_64-linux-gnu/perl-base/
export $PERL5LIB
sh ./cuda_8.0.61_375.26_linux.run --override

sudo rm /usr/lib/x86_64-linux-gnu/perl-base/InstallUtils.pm

Durante l'installazione non si deve installare il driver e si deve accettare l'installazione su piattaforma non supportata

E' conveniente installare anche gli esempi.