To pytanie ma już odpowiedź tutaj:

Wdrażam pewien algorytm przetwarzania obrazu w CUDA i mam kilka pytań dotyczących ogólnego problemu synchronizacji wątków.

Przedstawiony problem można wyjaśnić w następujący sposób:

Mamy obraz o rozmiarze W * H. Dla każdego piksela obrazu muszę uruchomić 9 równoległych procesów o identycznych danych, a każdy proces daje w wyniku tablicę wartości (tablice mają tę samą długość dla całego algorytmu, powiedzmy N, czyli około 20 lub 30 ). Dla każdego piksela te 9 procesów gromadzi swoje wyniki w końcowej tablicy (pojedyncza tablica dla każdego piksela) po zakończeniu obliczeń.

Aby to zrównoważyć, zaprojektowałem następującą strukturę: generuję bloki o wymiarach (10,10,9), co oznacza, że każdy blok wątku przetworzy obraz podrzędny o rozmiarze 10 * 10, a każdy wątek przetworzy 1 z 9 identycznych procesów dla jednego piksela. Wymiar siatki będzie w tym przypadku (W / 10, H / 10,1). Dla bloku wątków przydzielę tablicę pamięci współużytkowanej o długości 100 * N, a każdy wątek zapisze w odpowiednim miejscu pamięci współdzielonej zgodnie ze współrzędnymi jego bieżącego piksela. Potrzebuję więc synchronizacji z atomicAdd i __synchthreads () tutaj.

Problem polega na tym, że jeśli piksel ma wartość zero, to nie musimy go wcale przetwarzać, więc chcę wyjść dla takich pikseli, w przeciwnym razie wykonam niepotrzebną pracę, ponieważ duża część obrazu składa się zer (tło). Pomyślałem więc o napisaniu czegoś takiego:

//X and Y are the coordinates of the current pixel in the input image.
//threadIdx.z gives the index of the process among the 9 for the current pixel. 

int X=blockIdx.x * blockDim.x + threadIdx.x;
int Y=blockIdx.y * blockDim.y + threadIdx.y;
int numOfProcessForTheCurrPixel=threadIdx.z;
int linearIndexOfPixelInBlock=threadIdx.y * blockDim.x + threadIdx.x;

unsigned short pixelValue=tex2D(image,X,Y);
//Here, threads processing zero-pixels will exit immediately.
if(pixelValue==0)
 return;

float resultArray[22];
//Fill the result array according to our algorithm, mostly irrelevant stuff.
ProcessPixel(resultArray,X,Y,numOfProcessForTheCurrPixel);

for(int i=0;i<22;i++)
    atomicAdd(&__sharedMemoryArray[22*linearIndexOfPixelInBlock + i],resultArray[i]);

 __syncthreads(); 
 //Then copy from the shared to the global memory and etc. 

W tej sytuacji martwi mnie to, co mówi Przewodnik po programowaniu:

__ syncthreads () jest dozwolony w kodzie warunkowym, ale tylko wtedy, gdy warunkowo ocenia identycznie w całym bloku wątku, w przeciwnym razie wykonanie kodu może się zawiesić lub wywołać niezamierzone skutki uboczn

Tak więc w moim przypadku, jeśli niektóre piksele w bloku wątków 10 * 10 mają zero, a niektóre lub nie, to wątki należące do zerowych pikseli zakończą się natychmiast na początku, a pozostałe wątki będą kontynuować przetwarzanie. Co z synchronizacją w tym przypadku, czy będzie ona nadal działać poprawnie, czy też będzie generować niezdefiniowane zachowanie, jak mówi Podręcznik programowania? Myślałem o tym, aby wątki o zerowym pikselu przetwarzały śmieci, aby były zajęte, ale niepotrzebnie zwiększy to czas przetwarzania, jeśli mamy bloki składające się całkowicie z zer (i mamy je bardzo często). Co można zrobić w tym przypadku?