Bir CUDA çekirdeğinin paylaşılan bellek ile <code>volatile</code> anahtar kelime kullanmanız gerekir hangi şartlar altında

Question

?

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

Ben bu durumda uçucu olması products gerekiyor mu: Ben volatile tüm değerleri önbelleğe asla derleyici söyler, ama benim sorum paylaşılan dizi davranışı hakkında olduğunu anlıyoruz? Sadece önbellek onu derleyici dizinin tamamı önbelleğe olabilir ve bu yüzden volatile olmaya gerek, ya olacak o mümkün mü herşey iplik 0. tarafından okunan sonunda hariç her dizi girişi yalnızca tek iş parçacığı tarafından erişilir elementler?

Teşekkürler!

Answer 1

Paylaşılan bir diziyi volatile olarak bildirmezseniz, derleyici, herhangi bir iş parçacığı için herhangi bir iş parçacığı için, yazmaçlarını (kapsamı tek bir iş parçacığına özgü olan) konumlandırarak paylaşılan hafızdaki konumları en iyi duruma getirmek için serbesttir. . Bu belirli paylaşımlı öğeye yalnızca bir iş parçacığından erişip erişmediğiniz doğrudur. Bu nedenle, bir bloğun iş parçacıkları arasında bir iletişim aracı olarak paylaşılan bellek kullanırsanız, bunu en iyi bildirmek en iyisidir volatile.

Açıkçası her iş parçacığı yalnızca paylaşılan bellek kendi unsurlarını erişilen eğer, ve asla başka bir iş parçacığı ile ilişkili olanlar, o zaman bu önemli değil, ve derleyici optimizasyon şey kırmak olmaz.

Durumunuzda, her bir iş parçacığının eriştiği bir kod bölümünüzün kendi paylaşılan bellek öğelerine sahip olduğu ve iyi anlaşılan bir konumda yalnızca iç içe geçmiş erişim gerçekleştiği durumda, derleyiciyi zorlamak için memory fence function kullanabilirsiniz. Kayıtlar içinde geçici olarak saklanan değerleri boşaltmak için paylaşılan diziye geri dönün. Yani __threadfence_block() yararlı olabileceğini düşünüyorum, ama senin durumunda, __syncthreads()already has memory-fencing functionality built in içinde olabilir. Dolayısıyla, __syncthreads() aramanız, iş parçacığı eşitlemeyi zorlamak için olduğu gibi, paylaşılan belleğe kaydedilmiş önbelleğe alınmış değerlerin paylaşılan belleğe geri gönderilmesini zorlamak için yeterlidir. Kodunuzun sonunda bu azalma performans endişe ise

arada, bunu hızlandırmak için bir paralel azaltma yöntemini kullanmayı düşünebilirsiniz.