ustalenie, ile bloków i wątków dla jądra cuda, i jak z nich korzystać

Question

Próbowałem dowiedzieć się, jak zrobić to, co uważałem za proste jądro, aby wziąć średnią z wartości w matrycy 2d, ale mam pewne problemy z procesem myślowym.

Zgodnie z danymi mojego urządzenia deviceQuery, mój układ GPU ma 16MP, 32cores/mp, bloki max to 1024x1024x64 i mam max wątków/bloków = 1024.

Pracuję więc nad przetwarzaniem niektórych dużych obrazów. Może 5000px x 3500px lub coś w tym stylu. Jedno z moich ziaren pobiera średnią wartości z wszystkich pikseli w obrazie.

Istniejący kod zawiera obrazy przechowywane jako tablica 2D [wiersze] [cols]. Tak, że jądro w C wygląda tak, jak można by oczekiwać, z pętlą nad wierszami i pętlą nad colami, z kalkulacją w środku.

Jak skonfigurować część obliczeniową wymiaru tego kodu w CUDA? Przyjrzałem się kodowi redukcji w SDK, ale dotyczy to tablicy jednowymiarowej. To nie ma żadnej wzmianki o tym, jak ustawić liczbę bloków i wątków, gdy masz tok 2D.

myślę ja faktycznie trzeba go ustawić tak jak i to, gdzie chciałbym, aby ktoś w gong i pomoc:

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

Czy to wydaje się mieć sens dla konfiguracji ?

Następnie do jądra i działa w danym rzędzie lub kolumnie, to muszę używać

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

Przynajmniej sądzę, że to działałoby na uzyskanie wiersza i kolumny.

Jak uzyskać dostęp do tego wiersza r i kolumny c w jądrze? W instrukcji programowania CUDA znalazłem następujący kod:

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

który wygląda podobnie do tego, jak byłoby użyć malloc w C zadeklarować tablicy 2D, ale to nie robi mieć żadnej wzmianki o dostępie tę tablicę w swoim jądrze . Zgaduję, że w moim kodzie użyję tego połączenia cudaMallocPitch, a następnie wykonam memcpy, aby wprowadzić moje dane do tablicy 2D na urządzeniu?

Wszelkie wskazówki są mile widziane! Dzięki!

Answer 1

Poniżej znajduje się krótki fragment z prostym kernelem z mojego własnego kodu. Wskaźniki pływakowe są wskaźnikami wszystkich urządzeń. Mam nadzieję, że to jest pomocne.

definiuje i pomoc funkcje:

#define BLOCK_SIZE 16 

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

blok obliczeniowy rozmiar:

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

wezwanie 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

W przypadku aplikacji wydajnościowych, takich jak ta, należy przechowywać informacje o macierzy 2D jako pojedynczą macierz w pamięci. Więc jeśli masz macierz M x N, możesz przechowywać ją w pojedynczej macierzy o długości M * N.

Więc jeśli chcesz zapisać macierz 2x2

(1 , 2) 
(3 , 4) 

Następnie należy utworzyć pojedynczą Array zainicjować elementy w wierszu I i stosując następujące kolumny j.

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; 

ten sposób biorąc 2D tablicy i przechowywanie go tam pojedynczy ciągły element pamięci w ten sposób nazywa się przechowywanie danych w rzędzie major kolejności. Zobacz artykuł z Wikipedii here. Po zmianie układu danych na format tego rodzaju można skorzystać z redukcji, która została pokazana w SDK, a kod powinien być dużo szybszy, ponieważ będzie można zrobić więcej scalonych odczytów w kodzie jądra GPU.

Answer 3

Niedawno wymyśliłem to pytanie w następujący sposób.

// 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 = numer bloku
blocksize = wątków na blok

Oto analogicznym 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); 
} 

Powodzenia !!!