Kiedy używać volatile ze wspólną pamięcią CUDA

Question

W jakich okolicznościach należy użyć słowa kluczowego volatile z pamięcią wspólną jądra CUDA? Rozumiem, że volatile informuje kompilator nigdy buforować żadnych wartości, ale moje pytanie jest o zachowanie ze wspólną tablicy:

__shared__ float products[THREADS_PER_ACTION]; 

// some computation 
products[threadIdx.x] = localSum; 

// wait for everyone to finish their computation 
__syncthreads(); 

// then a (basic, ugly) reduction: 
if (threadIdx.x == 0) { 
    float globalSum = 0.0f; 
    for (i = 0; i < THREADS_PER_ACTION; i++) 
     globalSum += products[i]; 
} 

muszę products być lotny w tym przypadku? Do każdego wpisu tablicy dostęp ma tylko jeden wątek, z wyjątkiem na końcu, gdzie wszystko jest czytane przez wątek 0. Czy możliwe jest, że kompilator może buforować całą tablicę, a więc potrzebuję go jako volatile, czy też będzie to tylko pamięć podręczna elementy?

Dzięki!

Answer 1

Jeśli nie zadeklarujesz udostępnionej tablicy jako volatile, wówczas kompilator może dowolnie optymalizować lokalizacje we wspólnej pamięci, umieszczając je w rejestrach (których zakres jest specyficzny dla pojedynczego wątku), dla dowolnego wątku, przy jego wyborze . Jest to prawdą, niezależnie od tego, czy masz dostęp do tego konkretnego elementu współdzielonego z tylko jednego wątku, czy nie. Dlatego jeśli korzystasz z pamięci współdzielonej jako nośnika komunikacji między wątkami bloku, najlepiej jest zadeklarować to jako volatile.

Oczywiście, jeśli każdy wątek ma dostęp tylko do własnych elementów pamięci współużytkowanej, a nigdy do tych powiązanych z innym wątkiem, nie ma to znaczenia, a optymalizacja kompilatora niczego nie złamie.

W twoim przypadku, gdy masz sekcję kodu, w której każdy wątek uzyskuje dostęp do własnych elementów pamięci współużytkowanej, a jedyny dostęp między wątkami odbywa się w dobrze zrozumiałej lokalizacji, możesz użyć memory fence function, aby wymusić kompilator eksmitować wszelkie wartości, które są tymczasowo przechowywane w rejestrach, z powrotem do wspólnej tablicy. Możesz więc pomyśleć, że __threadfence_block() może być przydatny, ale w twoim przypadku, __syncthreads()already has memory-fencing functionality built in. Zatem twoje wywołanie __syncthreads() jest wystarczające, aby wymusić synchronizację nici, jak również wymuszenie, aby wszystkie buforowane wartości z pamięci współdzielonej zostały wyrzucone z powrotem do pamięci współdzielonej.

Nawiasem mówiąc, jeśli to zmniejszenie na końcu kodu ma wpływ na wydajność, można rozważyć zastosowanie metody redukcji równoległej w celu przyspieszenia.