martedì 18 aprile 2023

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;
 }
 

 

Nessun commento:

Posta un commento

Debugger integrato ESP32S3

Aggiornamento In realta' il Jtag USB funziona anche sui moduli cinesi Il problema risiede  nell'ID USB della porta Jtag. Nel modulo...