determinare quanti blocchi e thread per un kernel CUDA e come usarli

Question

Ho cercato di capire come rendere quello che pensavo sarebbe un semplice kernel per prendere la media dei valori in una matrice 2d, ma sto riscontrando alcuni problemi che riguardano il mio processo di pensiero.

In base al mio output deviceQuery, la mia GPU ha 16MP, 32cores/mp, i blocchi max è 1024x1024x64 e ho un massimo thread/blocco = 1024.

Quindi, sto lavorando ad alcune immagini di grandi dimensioni. Forse 5000px x 3500px o qualcosa del genere. Uno dei miei kernel sta prendendo una media di alcuni valori su tutti i pixel dell'immagine.

Il codice esistente ha le immagini memorizzate come un array 2D [righe] [colonne]. In modo che il kernel, in C, sembri come ci si aspetterebbe, con un ciclo su righe e un ciclo su colonne, con il calcolo nel mezzo.

Quindi, come si imposta la porzione di calcolo della quota di questo codice in CUDA? Ho visto il codice di riduzione nell'SDK, ma quello è per un array a dimensione singola. Non ha alcuna menzione su come impostare il numero di blocchi e fili per quando si dispone di 2D.

Sto pensando che avrei davvero bisogno di configurarlo in questo modo, e questo è dove mi piacerebbe che qualcuno carillon e aiuto:

num_threads=1024; 
blocksX = num_cols/sqrt(num_threads); 
blocksY = num_rows/sqrt(num_threads); 
num_blocks = (num_rows*num_cols)/(blocksX*blocksY); 

dim3 dimBlock(blocksX, blocksY, 1); 
dim3 dimGrid(num_blocks, 1, 1); 

Questo sembra avere senso per la configurazione ?

E poi nel kernel, per lavorare su una particolare riga o colonna, avrei dovuto usare

rowidx = (blockIdx.x * blockDim.x) + threadId.x colidx = (blockIdx. y * blockDim.y) + threadId.y

Almeno penso che avrebbe funzionato per ottenere una riga e una colonna.

Come potrei quindi accedere a quella particolare riga r e colonna c nel kernel? Nella guida di programmazione CUDA ho trovato il seguente codice:

// Host code int width = 64, height = 64; 
float* devPtr; size_t pitch; 
cudaMallocPitch(&devPtr, &pitch, width * sizeof(float), height); 
MyKernel<<<100, 512>>>(devPtr, pitch, width, height); 
// Device code __global__ void MyKernel(float* devPtr, size_t pitch, int width, int height) 
{ 
for (int r = 0; r < height; ++r) 
{ 
float* row = (float*)((char*)devPtr + r * pitch); 
for (int c = 0; c < width; ++c) 
{ 
float element = row[c]; 
} 
} 
} 

che sembra simile a come devi usare malloc in C per dichiarare una matrice 2D, ma Esso non dispone alcuna menzione di accesso che matrice in un proprio kernel . Immagino nel mio codice, userò quella chiamata cudaMallocPitch e poi eseguirò una memcpy per ottenere i miei dati nell'array 2D sul dispositivo?

Qualsiasi consiglio apprezzato! Grazie!

Answer 1

Di seguito è riportato un breve snippet con un semplice kernel dal mio codice. I puntatori mobili sono tutti puntatori di dispositivo. Spero che questo sia utile.

definisce e aiuto funzioni:

#define BLOCK_SIZE 16 

int iDivUp(int a, int b){ 
    return (a % b != 0) ? (a/b + 1) : (a/b); 
} 

blocco di calcolo dimensioni:

dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE); 
dim3 dimGridProj(iDivUp(width,BLOCK_SIZE), iDivUp(height,BLOCK_SIZE)); 

chiamata Host:

calc_residual<<<dimGridProj, dimBlock>>>(d_image1, d_proj1, d_raynorm1, d_resid1, width, height); 

Kernel:

__global__ void calc_residual(float *d_imagep, float *d_projp, float *d_raysump, float *d_residualp, int width, int height) 
{ 
    int iy = blockDim.y * blockIdx.y + threadIdx.y; 
if (iy >= height) { 
    return; 
} 
int ix = blockDim.x * blockIdx.x + threadIdx.x; 
if (ix >= width) { 
    return; 
} 
int idx = iy * width + ix; 
float raysumv = d_raysump[idx]; 
if (raysumv > 0.001) { 
    d_residualp[idx] = (d_projp[idx]-d_imagep[idx])/raysumv; 
} 
else{ 
    d_residualp[idx] = 0; 
} 
} 

Answer 2

Per applicazioni di prestazioni come questa è necessario memorizzare le informazioni della matrice 2D come un singolo array in memoria. Quindi, se hai una matrice M x N, puoi conservarla in un unico array di lunghezza M * N.

Quindi, se si desidera memorizzare la matrice 2x2

(1 , 2) 
(3 , 4) 

quindi si crea una singola matrice di inizializzare gli elementi in riga i e j colonna utilizzando la seguente.

int rows=2; 
int cols=2; 
float* matrix = malloc(sizeof(float)*rows*cols); 
matrix[i*cols+j]=yourValue; 
//element 0,0 
matrix[0*cols+0]=1.0; 
//element 0,1 
matrix[0*cols+1]=2.0; 
//element 1,0 
matrix[1*cols+0]=3.0; 
//element 1,1 
matrix[1*cols+1]=4.0; 

Questo modo di prendere un array 2D e riporlo un unico pezzo continuo di memoria in questo modo viene chiamato memorizzazione dei dati al fine di riga principale. Vedi l'articolo di Wikipedia here. Una volta modificato il layout dei dati in questo tipo di formato, puoi utilizzare la riduzione mostrata nell'SDK e il tuo codice dovrebbe essere molto più veloce in quanto sarai in grado di eseguire più letture a coalescenza nel codice del kernel GPU.

Answer 3

Recentemente, ho trovato questa domanda nella seguente maniera.

// Grid and block size 
const dim3 blockSize(16,16,1); 
const dim3 gridSize(numRows, numCols, 1); 
// kernel call 
rgba_to_greyscale<<<gridSize, blockSize>>>(d_rgbaImage, d_greyImage, numRows, numCols 

gridSize = Numero di blocco
blocksize = filetti per Block

Ecco corrispondente kernel

__global__ void rgba_to_greyscale(const uchar4* const rgbaImage, 
         unsigned char* const greyImage, 
         int numRows, int numCols) 
{ 
    int idx = blockIdx.x + blockIdx.y * numRows; 
    uchar4 pixel  = rgbaImage[idx]; 
    float intensity = 0.299f * pixel.x + 0.587f * pixel.y + 0.114f * pixel.z; 
    greyImage[idx] = static_cast<unsigned char>(intensity); 
} 

Buona fortuna !!!