CUDA. Jak rozwinąć pierwsze 32 wątki, aby były wykonywane równolegle?

Question

Wiem, że "każda osnowa zawiera wątki kolejnych, rosnących identyfikatorów wątków z pierwszą osnową zawierającą wątek 0", więc pierwsze 32 wątki powinny znajdować się w pierwszej osnowie. Ponadto wiem, że wszystkie wątki w jednej osnowie są wykonywane jednocześnie na dowolnym dostępnym multiprocesorze strumieniowym.

Jak zrozumiałem, z tego powodu nie ma potrzeby synchronizowania wątku, jeśli wykonywana jest tylko jedna osnowa. Jednak poniższy kod generuje błędną odpowiedź, jeśli usunę dowolny z __syncthreads() w przedostatnim bloku if. Próbowałem znaleźć przyczynę, ale skończyłem z niczym. Naprawdę mam nadzieję na twoją pomoc, więc możesz mi powiedzieć, co jest nie tak z tym kodem? Dlaczego nie mogę opuścić tylko ostatniego __syncthreads() i uzyskać prawidłową odpowiedź?

#define BLOCK_SIZE 128 

__global__ void reduce (int * inData, int * outData) 
{ 
__shared__ int data [BLOCK_SIZE]; 
int tid = threadIdx.x; 
int i = blockIdx.x * blockDim.x + threadIdx.x; 

data [tid] = inData [i] + inData [i + blockDim.x/2 ]; 
__syncthreads(); 

for (int s = blockDim.x/4; s > 32; s >>= 1) 
{ 
    if (tid < s) 
    data [tid] += data [tid + s]; 
    __syncthreads(); 
} 

if (tid < 32) 
{ 
    data [tid] += data [tid + 32]; 
    __syncthreads(); 
    data [tid] += data [tid + 16]; 
    __syncthreads(); 
    data [tid] += data [tid + 8]; 
    __syncthreads(); 
    data [tid] += data [tid + 4]; 
    __syncthreads(); 
    data [tid] += data [tid + 2]; 
    __syncthreads(); 
    data [tid] += data [tid + 1]; 
    __syncthreads(); 
} 
if (tid == 0) 
    outData [blockIdx.x] = data [0]; 
} 

void main() 
{ 
... 
reduce<<<dim3(128), dim3(128)>>>(dev_data, dev_res); 
... 
} 

P.S. Używam GT560Ti

Answer 1

Należy zadeklarować zmienną pamięci dzielonej jako lotny:

__shared__ volatile int data [BLOCK_SIZE]; 

Problem widzisz jest artefaktem architekturze Fermi i optymalizacji kompilatora. W architekturze Fermi brakuje instrukcji do obsługi bezpośrednio w pamięci współdzielonej (były one obecne w serii G80/90/GT200). Tak więc wszystko jest ładowane, aby zarejestrować się, zmanipulować i zapisać do pamięci współdzielonej. Jednak kompilator może wywnioskować, że kod mógłby zostać przyspieszony, gdyby szereg operacji został wykonany w rejestrze, bez pośrednich obciążeń i zapisów z/do pamięci współdzielonej. Jest to całkowicie w porządku, z wyjątkiem z wyjątkiem, gdy polegasz na niejawnej synchronizacji wątków w obrębie tej samej osnowy manipulującej wspólną pamięcią, jak w tym rodzaju kodu redukcji.

Deklarując, że bufor pamięci współużytkowanej jest niestabilny, zmuszasz kompilator do wymuszania zapisu pamięci współdzielonej po każdym etapie redukcji i przywracana jest niejawna synchronizacja danych między wątkami w warp.

Kwestia ta została omówiona w uwagach do programowania dla Fermiego, który jest wysyłany (lub może być dostarczany) za pomocą zestawu narzędzi CUDA.